fix: lintrunner --all-files --take NEWLINE -a-pytorch-AtomGit

MCMakeLists.txt+396-397

MThird_Party_Open_Source_Software_Notice+190-190

Mrequirements.txt+7-7

	@@ -1,7 +1,7 @@
1	--extra-index-url https://download.pytorch.org/whl/cpu	1	--extra-index-url https://download.pytorch.org/whl/cpu
2		2
3	pyyaml	3	pyyaml
4	setuptools	4	setuptools
5	auditwheel	5	auditwheel
6		6
7	torch==2.7.1	7	torch==2.7.1

Msetup.py+782-782

	@@ -1,782 +1,782 @@
1	import glob	1	import glob
2	import multiprocessing	2	import multiprocessing
3	import multiprocessing.pool	3	import multiprocessing.pool
4	import os	4	import os
5	import re	5	import re
6	import shutil	6	import shutil
7	import stat	7	import stat
8	import subprocess	8	import subprocess
9	import sys	9	import sys
10	import traceback	10	import traceback
11	import platform	11	import platform
12	import time	12	import time
13	import sysconfig	13	import sysconfig
14	from sysconfig import get_paths	14	from sysconfig import get_paths
15	from pathlib import Path	15	from pathlib import Path
16	from typing import Union	16	from typing import Union
17		17
18	import distutils.ccompiler	18	import distutils.ccompiler
19	import distutils.command.clean	19	import distutils.command.clean
20	from distutils.version import LooseVersion	20	from distutils.version import LooseVersion
21	from distutils.command.build_py import build_py	21	from distutils.command.build_py import build_py
22	from setuptools import setup, distutils, Extension, find_packages	22	from setuptools import setup, distutils, Extension, find_packages
23	from setuptools.command.build_clib import build_clib	23	from setuptools.command.build_clib import build_clib
24	from setuptools.command.build_ext import build_ext	24	from setuptools.command.build_ext import build_ext
25	from setuptools.command.egg_info import egg_info	25	from setuptools.command.egg_info import egg_info
26	from setuptools.command.install import install	26	from setuptools.command.install import install
27	from wheel.bdist_wheel import bdist_wheel	27	from wheel.bdist_wheel import bdist_wheel
28		28
29	# Disable autoloading before running 'import torch' to avoid circular dependencies	29	# Disable autoloading before running 'import torch' to avoid circular dependencies
30	os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"	30	os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"
31		31
32	from torchnpugen.utils import PathManager	32	from torchnpugen.utils import PathManager
33		33
34	BASE_DIR = os.path.dirname(os.path.realpath(__file__))	34	BASE_DIR = os.path.dirname(os.path.realpath(__file__))
35	THIRD_PARTY_PATH = os.path.join(BASE_DIR, "third_party")	35	THIRD_PARTY_PATH = os.path.join(BASE_DIR, "third_party")
36	PathManager.check_directory_path_readable(os.path.join(BASE_DIR, "version.txt"))	36	PathManager.check_directory_path_readable(os.path.join(BASE_DIR, "version.txt"))
37	with open(os.path.join(BASE_DIR, "version.txt")) as version_f:	37	with open(os.path.join(BASE_DIR, "version.txt")) as version_f:
38	VERSION = version_f.read().strip()	38	VERSION = version_f.read().strip()
39	UNKNOWN = "Unknown"	39	UNKNOWN = "Unknown"
40	BUILD_PERMISSION = stat.S_IRUSR \| stat.S_IWUSR \| stat.S_IXUSR \| stat.S_IRGRP \| stat.S_IXGRP	40	BUILD_PERMISSION = stat.S_IRUSR \| stat.S_IWUSR \| stat.S_IXUSR \| stat.S_IRGRP \| stat.S_IXGRP
41		41
42	DISABLE_TORCHAIR = "FALSE"	42	DISABLE_TORCHAIR = "FALSE"
43	if os.environ.get("DISABLE_INSTALL_TORCHAIR") is not None:	43	if os.environ.get("DISABLE_INSTALL_TORCHAIR") is not None:
44	DISABLE_TORCHAIR = os.environ.get("DISABLE_INSTALL_TORCHAIR")	44	DISABLE_TORCHAIR = os.environ.get("DISABLE_INSTALL_TORCHAIR")
45	DISABLE_RPC = "FALSE"	45	DISABLE_RPC = "FALSE"
46	if os.environ.get("DISABLE_RPC_FRAMEWORK") is not None:	46	if os.environ.get("DISABLE_RPC_FRAMEWORK") is not None:
47	DISABLE_RPC = os.environ.get("DISABLE_RPC_FRAMEWORK")	47	DISABLE_RPC = os.environ.get("DISABLE_RPC_FRAMEWORK")
48	ENABLE_LTO = "FALSE"	48	ENABLE_LTO = "FALSE"
49	if os.environ.get("ENABLE_LTO") is not None:	49	if os.environ.get("ENABLE_LTO") is not None:
50	ENABLE_LTO = os.environ.get("ENABLE_LTO")	50	ENABLE_LTO = os.environ.get("ENABLE_LTO")
51	PGO_MODE = 0	51	PGO_MODE = 0
52	if os.environ.get("PGO_MODE") is not None:	52	if os.environ.get("PGO_MODE") is not None:
53	PGO_MODE = int(os.environ.get("PGO_MODE"))	53	PGO_MODE = int(os.environ.get("PGO_MODE"))
54		54
55	# change to use cxx11.abi in default since 2.7	55	# change to use cxx11.abi in default since 2.7
56	USE_CXX11_ABI = True	56	USE_CXX11_ABI = True
57	if os.environ.get("_GLIBCXX_USE_CXX11_ABI") is not None and os.environ.get("_GLIBCXX_USE_CXX11_ABI") == "0":	57	if os.environ.get("_GLIBCXX_USE_CXX11_ABI") is not None and os.environ.get("_GLIBCXX_USE_CXX11_ABI") == "0":
58	USE_CXX11_ABI = False	58	USE_CXX11_ABI = False
59		59
60		60
61	def get_submodule_folders():	61	def get_submodule_folders():
62	git_modules_path = os.path.join(BASE_DIR, ".gitmodules")	62	git_modules_path = os.path.join(BASE_DIR, ".gitmodules")
63	default_modules_path = [	63	default_modules_path = [
64	os.path.join(THIRD_PARTY_PATH, name)	64	os.path.join(THIRD_PARTY_PATH, name)
65	for name in [	65	for name in [
66	"op-plugin",	66	"op-plugin",
67	]	67	]
68	]	68	]
69	if not os.path.exists(git_modules_path):	69	if not os.path.exists(git_modules_path):
70	return default_modules_path	70	return default_modules_path
71	with open(git_modules_path) as f:	71	with open(git_modules_path) as f:
72	return [	72	return [
73	os.path.join(BASE_DIR, line.split("=", 1)[1].strip())	73	os.path.join(BASE_DIR, line.split("=", 1)[1].strip())
74	for line in f.readlines()	74	for line in f.readlines()
75	if line.strip().startswith("path")	75	if line.strip().startswith("path")
76	]	76	]
77		77
78		78
79	def check_submodules():	79	def check_submodules():
80	def not_exists_or_empty(folder):	80	def not_exists_or_empty(folder):
81	return not os.path.exists(folder) or (	81	return not os.path.exists(folder) or (
82	os.path.isdir(folder) and len(os.listdir(folder)) == 0	82	os.path.isdir(folder) and len(os.listdir(folder)) == 0
83	)	83	)
84		84
85	folders = get_submodule_folders()	85	folders = get_submodule_folders()
86	# If none of the submodule folders exists, try to initialize them	86	# If none of the submodule folders exists, try to initialize them
87	if all(not_exists_or_empty(folder) for folder in folders):	87	if all(not_exists_or_empty(folder) for folder in folders):
88	try:	88	try:
89	print(" --- Trying to initialize submodules")	89	print(" --- Trying to initialize submodules")
90	start = time.time()	90	start = time.time()
91	subprocess.check_call(["git", "submodule", "update", "--init", "--recursive"], cwd=BASE_DIR) # Compliant	91	subprocess.check_call(["git", "submodule", "update", "--init", "--recursive"], cwd=BASE_DIR) # Compliant
92	end = time.time()	92	end = time.time()
93	print(f" --- Submodule initialization took {end - start:.2f} sec")	93	print(f" --- Submodule initialization took {end - start:.2f} sec")
94	except Exception:	94	except Exception:
95	print(" --- Submodule initalization failed")	95	print(" --- Submodule initalization failed")
96	print("Please run:\n\tgit submodule init && git submodule update")	96	print("Please run:\n\tgit submodule init && git submodule update")
97	sys.exit(1)	97	sys.exit(1)
98		98
99		99
100	check_submodules()	100	check_submodules()
101		101
102		102
103	def get_sha(pytorch_root: Union[str, Path]) -> str:	103	def get_sha(pytorch_root: Union[str, Path]) -> str:
104	try:	104	try:
105	return (	105	return (
106	subprocess.check_output(["git", "rev-parse", "HEAD"], cwd=pytorch_root) # Compliant	106	subprocess.check_output(["git", "rev-parse", "HEAD"], cwd=pytorch_root) # Compliant
107	.decode("ascii")	107	.decode("ascii")
108	.strip()	108	.strip()
109	)	109	)
110	except Exception:	110	except Exception:
111	return UNKNOWN	111	return UNKNOWN
112		112
113		113
114	def generate_torch_npu_version():	114	def generate_torch_npu_version():
115	torch_npu_root = Path(__file__).parent	115	torch_npu_root = Path(__file__).parent
116	version_path = torch_npu_root / "torch_npu" / "version.py"	116	version_path = torch_npu_root / "torch_npu" / "version.py"
117	if version_path.exists():	117	if version_path.exists():
118	version_path.unlink()	118	version_path.unlink()
119	flags = os.O_WRONLY \| os.O_CREAT \| os.O_EXCL	119	flags = os.O_WRONLY \| os.O_CREAT \| os.O_EXCL
120	modes = stat.S_IWUSR \| stat.S_IRUSR	120	modes = stat.S_IWUSR \| stat.S_IRUSR
121	sha = get_sha(torch_npu_root)	121	sha = get_sha(torch_npu_root)
122	if os.getenv("BUILD_WITHOUT_SHA") is None:	122	if os.getenv("BUILD_WITHOUT_SHA") is None:
123	global VERSION	123	global VERSION
124	VERSION += "+git" + sha[:7]	124	VERSION += "+git" + sha[:7]
125	with os.fdopen(os.open(version_path, flags, modes), 'w') as f:	125	with os.fdopen(os.open(version_path, flags, modes), 'w') as f:
126	f.write("__version__ = '{version}'\n".format(version=VERSION))	126	f.write("__version__ = '{version}'\n".format(version=VERSION))
127	f.write("git_version = {}\n".format(repr(sha)))	127	f.write("git_version = {}\n".format(repr(sha)))
128	os.chmod(version_path, mode=stat.S_IRUSR \| stat.S_IEXEC \| stat.S_IRGRP \| stat.S_IXGRP)	128	os.chmod(version_path, mode=stat.S_IRUSR \| stat.S_IEXEC \| stat.S_IRGRP \| stat.S_IXGRP)
129		129
130		130
131	generate_torch_npu_version()	131	generate_torch_npu_version()
132		132
133		133
134	def which(thefile):	134	def which(thefile):
135	path = os.environ.get("PATH", os.defpath).split(os.pathsep)	135	path = os.environ.get("PATH", os.defpath).split(os.pathsep)
136	for d in path:	136	for d in path:
137	fname = os.path.join(d, thefile)	137	fname = os.path.join(d, thefile)
138	fnames = [fname]	138	fnames = [fname]
139	if sys.platform == 'win32':	139	if sys.platform == 'win32':
140	exts = os.environ.get('PATHEXT', '').split(os.pathsep)	140	exts = os.environ.get('PATHEXT', '').split(os.pathsep)
141	fnames += [fname + ext for ext in exts]	141	fnames += [fname + ext for ext in exts]
142	for name in fnames:	142	for name in fnames:
143	if os.access(name, os.F_OK \| os.X_OK) and not os.path.isdir(name):	143	if os.access(name, os.F_OK \| os.X_OK) and not os.path.isdir(name):
144	return name	144	return name
145	return None	145	return None
146		146
147		147
148	def get_cmake_command():	148	def get_cmake_command():
149	def _get_version(cmd):	149	def _get_version(cmd):
150	for line in subprocess.check_output([cmd, '--version']).decode('utf-8').split('\n'):	150	for line in subprocess.check_output([cmd, '--version']).decode('utf-8').split('\n'):
151	if 'version' in line:	151	if 'version' in line:
152	return LooseVersion(line.strip().split(' ')[2])	152	return LooseVersion(line.strip().split(' ')[2])
153	raise RuntimeError('no version found')	153	raise RuntimeError('no version found')
154	"Returns cmake command."	154	"Returns cmake command."
155	cmake_command = 'cmake'	155	cmake_command = 'cmake'
156	if platform.system() == 'Windows':	156	if platform.system() == 'Windows':
157	return cmake_command	157	return cmake_command
158	cmake3 = which('cmake3')	158	cmake3 = which('cmake3')
159	cmake = which('cmake')	159	cmake = which('cmake')
160	if cmake3 is not None and _get_version(cmake3) >= LooseVersion("3.18.0"):	160	if cmake3 is not None and _get_version(cmake3) >= LooseVersion("3.18.0"):
161	cmake_command = 'cmake3'	161	cmake_command = 'cmake3'
162	return cmake_command	162	return cmake_command
163	elif cmake is not None and _get_version(cmake) >= LooseVersion("3.18.0"):	163	elif cmake is not None and _get_version(cmake) >= LooseVersion("3.18.0"):
164	return cmake_command	164	return cmake_command
165	else:	165	else:
166	raise RuntimeError('no cmake or cmake3 with version >= 3.18.0 found')	166	raise RuntimeError('no cmake or cmake3 with version >= 3.18.0 found')
167		167
168		168
169	def get_build_type():	169	def get_build_type():
170	build_type = 'Release'	170	build_type = 'Release'
171	if os.getenv('DEBUG', default='0').upper() in ['ON', '1', 'YES', 'TRUE', 'Y']:	171	if os.getenv('DEBUG', default='0').upper() in ['ON', '1', 'YES', 'TRUE', 'Y']:
172	build_type = 'Debug'	172	build_type = 'Debug'
173		173
174	if os.getenv('REL_WITH_DEB_INFO', default='0').upper() in ['ON', '1', 'YES', 'TRUE', 'Y']:	174	if os.getenv('REL_WITH_DEB_INFO', default='0').upper() in ['ON', '1', 'YES', 'TRUE', 'Y']:
175	build_type = 'RelWithDebInfo'	175	build_type = 'RelWithDebInfo'
176		176
177	return build_type	177	return build_type
178		178
179		179
180	def _get_build_mode():	180	def _get_build_mode():
181	for i in range(1, len(sys.argv)):	181	for i in range(1, len(sys.argv)):
182	if not sys.argv[i].startswith('-'):	182	if not sys.argv[i].startswith('-'):
183	return sys.argv[i]	183	return sys.argv[i]
184		184
185	raise RuntimeError("Run setup.py without build mode.")	185	raise RuntimeError("Run setup.py without build mode.")
186		186
187		187
188	def get_pytorch_dir():	188	def get_pytorch_dir():
189	try:	189	try:
190	import torch	190	import torch
191	return os.path.dirname(os.path.realpath(torch.__file__))	191	return os.path.dirname(os.path.realpath(torch.__file__))
192	except Exception:	192	except Exception:
193	_, _, exc_traceback = sys.exc_info()	193	_, _, exc_traceback = sys.exc_info()
194	frame_summary = traceback.extract_tb(exc_traceback)[-1]	194	frame_summary = traceback.extract_tb(exc_traceback)[-1]
195	return os.path.dirname(frame_summary.filename)	195	return os.path.dirname(frame_summary.filename)
196		196
197		197
198	def generate_bindings_code(base_dir):	198	def generate_bindings_code(base_dir):
199	python_execute = sys.executable	199	python_execute = sys.executable
200	generate_code_cmd = ["bash", os.path.join(base_dir, 'generate_code.sh'), python_execute, VERSION]	200	generate_code_cmd = ["bash", os.path.join(base_dir, 'generate_code.sh'), python_execute, VERSION]
201	if subprocess.call(generate_code_cmd) != 0: # Compliant	201	if subprocess.call(generate_code_cmd) != 0: # Compliant
202	print(	202	print(
203	'Failed to generate ATEN bindings: {}'.format(generate_code_cmd),	203	'Failed to generate ATEN bindings: {}'.format(generate_code_cmd),
204	file=sys.stderr)	204	file=sys.stderr)
205	sys.exit(1)	205	sys.exit(1)
206		206
207		207
208	def build_stub(base_dir):	208	def build_stub(base_dir):
209	build_stub_cmd = ["sh", os.path.join(base_dir, 'third_party/acl/libs/build_stub.sh')]	209	build_stub_cmd = ["sh", os.path.join(base_dir, 'third_party/acl/libs/build_stub.sh')]
210	if subprocess.call(build_stub_cmd) != 0:	210	if subprocess.call(build_stub_cmd) != 0:
211	print(	211	print(
212	'Failed to build stub: {}'.format(build_stub_cmd),	212	'Failed to build stub: {}'.format(build_stub_cmd),
213	file=sys.stderr)	213	file=sys.stderr)
214	sys.exit(1)	214	sys.exit(1)
215		215
216		216
217	def check_torchair_valid(base_dir):	217	def check_torchair_valid(base_dir):
218	# build with submodule of torchair, if path of torchair is valid	218	# build with submodule of torchair, if path of torchair is valid
219	torchair_path = os.path.join(base_dir, 'third_party/torchair/torchair')	219	torchair_path = os.path.join(base_dir, 'third_party/torchair/torchair')
220	return os.path.exists(torchair_path) and (	220	return os.path.exists(torchair_path) and (
221	os.path.isdir(torchair_path) and len(os.listdir(torchair_path)) != 0	221	os.path.isdir(torchair_path) and len(os.listdir(torchair_path)) != 0
222	)	222	)
223		223
224		224
225	def check_tensorpipe_valid(base_dir):	225	def check_tensorpipe_valid(base_dir):
226	tensorpipe_path = os.path.join(base_dir, 'third_party/Tensorpipe/tensorpipe')	226	tensorpipe_path = os.path.join(base_dir, 'third_party/Tensorpipe/tensorpipe')
227	return os.path.exists(tensorpipe_path)	227	return os.path.exists(tensorpipe_path)
228		228
229		229
230	def generate_dbg_files_and_strip():	230	def generate_dbg_files_and_strip():
231	library_dir = Path(BASE_DIR).joinpath("build/packages/torch_npu")	231	library_dir = Path(BASE_DIR).joinpath("build/packages/torch_npu")
232	dbg_dir = Path(BASE_DIR).joinpath("build/dbg")	232	dbg_dir = Path(BASE_DIR).joinpath("build/dbg")
233	os.makedirs(dbg_dir, exist_ok=True)	233	os.makedirs(dbg_dir, exist_ok=True)
234	library_files = [Path(i) for i in library_dir.rglob('*.so')]	234	library_files = [Path(i) for i in library_dir.rglob('*.so')]
235	for library_file in library_files:	235	for library_file in library_files:
236	subprocess.check_call(["eu-strip", library_file, "-f",	236	subprocess.check_call(["eu-strip", library_file, "-f",
237	str(dbg_dir.joinpath(library_file.name)) + ".debug"], cwd=BASE_DIR) # Compliant	237	str(dbg_dir.joinpath(library_file.name)) + ".debug"], cwd=BASE_DIR) # Compliant
238		238
239		239
240	def patchelf_dynamic_library():	240	def patchelf_dynamic_library():
241	# Process all .so files in lib directory	241	# Process all .so files in lib directory
242	lib_dir = Path(BASE_DIR).joinpath("build/packages/torch_npu/lib")	242	lib_dir = Path(BASE_DIR).joinpath("build/packages/torch_npu/lib")
243	lib_files = [str(i) for i in lib_dir.rglob('*.so')]	243	lib_files = [str(i) for i in lib_dir.rglob('*.so')]
244		244
245	for library_file in lib_files:	245	for library_file in lib_files:
246	subprocess.check_call(["patchelf", "--remove-needed", "libgomp.so.1", library_file], cwd=BASE_DIR) # Compliant	246	subprocess.check_call(["patchelf", "--remove-needed", "libgomp.so.1", library_file], cwd=BASE_DIR) # Compliant
247		247
248		248
249	def CppExtension(name, sources, args, *kwargs):	249	def CppExtension(name, sources, args, *kwargs):
250	r'''	250	r'''
251	Creates a :class:`setuptools.Extension` for C++.	251	Creates a :class:`setuptools.Extension` for C++.
252	'''	252	'''
253	pytorch_dir = get_pytorch_dir()	253	pytorch_dir = get_pytorch_dir()
254	temp_include_dirs = kwargs.get('include_dirs', [])	254	temp_include_dirs = kwargs.get('include_dirs', [])
255	temp_include_dirs.append(os.path.join(pytorch_dir, 'include'))	255	temp_include_dirs.append(os.path.join(pytorch_dir, 'include'))
256	temp_include_dirs.append(os.path.join(pytorch_dir, 'include/torch/csrc/api/include'))	256	temp_include_dirs.append(os.path.join(pytorch_dir, 'include/torch/csrc/api/include'))
257	kwargs['include_dirs'] = temp_include_dirs	257	kwargs['include_dirs'] = temp_include_dirs
258		258
259	temp_library_dirs = kwargs.get('library_dirs', [])	259	temp_library_dirs = kwargs.get('library_dirs', [])
260	temp_library_dirs.append(os.path.join(pytorch_dir, 'lib'))	260	temp_library_dirs.append(os.path.join(pytorch_dir, 'lib'))
261	temp_library_dirs.append(os.path.join(BASE_DIR, "third_party/acl/libs"))	261	temp_library_dirs.append(os.path.join(BASE_DIR, "third_party/acl/libs"))
262	kwargs['library_dirs'] = temp_library_dirs	262	kwargs['library_dirs'] = temp_library_dirs
263		263
264	libraries = kwargs.get('libraries', [])	264	libraries = kwargs.get('libraries', [])
265	libraries.append('c10')	265	libraries.append('c10')
266	libraries.append('torch')	266	libraries.append('torch')
267	libraries.append('torch_cpu')	267	libraries.append('torch_cpu')
268	libraries.append('torch_python')	268	libraries.append('torch_python')
269	libraries.append('hccl')	269	libraries.append('hccl')
270	kwargs['libraries'] = libraries	270	kwargs['libraries'] = libraries
271	kwargs['language'] = 'c++'	271	kwargs['language'] = 'c++'
272	return Extension(name, sources, args, *kwargs)	272	return Extension(name, sources, args, *kwargs)
273		273
274		274
275	class Clean(distutils.command.clean.clean):	275	class Clean(distutils.command.clean.clean):
276		276
277	def run(self):	277	def run(self):
278	f_ignore = open('.gitignore', 'r')	278	f_ignore = open('.gitignore', 'r')
279	ignores = f_ignore.read()	279	ignores = f_ignore.read()
280	pat = re.compile(r'^#( BEGIN NOT-CLEAN-FILES )?')	280	pat = re.compile(r'^#( BEGIN NOT-CLEAN-FILES )?')
281	for wildcard in filter(None, ignores.split('\n')):	281	for wildcard in filter(None, ignores.split('\n')):
282	match = pat.match(wildcard)	282	match = pat.match(wildcard)
283	if match:	283	if match:
284	if match.group(1):	284	if match.group(1):
285	# Marker is found and stop reading .gitignore.	285	# Marker is found and stop reading .gitignore.
286	break	286	break
287	# Ignore lines which begin with '#'.	287	# Ignore lines which begin with '#'.
288	else:	288	else:
289	for filename in glob.glob(wildcard):	289	for filename in glob.glob(wildcard):
290	if os.path.islink(filename):	290	if os.path.islink(filename):
291	raise RuntimeError(f"Failed to remove path: {filename}")	291	raise RuntimeError(f"Failed to remove path: {filename}")
292	if os.path.exists(filename):	292	if os.path.exists(filename):
293	try:	293	try:
294	shutil.rmtree(filename, ignore_errors=True)	294	shutil.rmtree(filename, ignore_errors=True)
295	except Exception as err:	295	except Exception as err:
296	raise RuntimeError(f"Failed to remove path: {filename}") from err	296	raise RuntimeError(f"Failed to remove path: {filename}") from err
297	f_ignore.close()	297	f_ignore.close()
298		298
299	# It's an old-style class in Python 2.7...	299	# It's an old-style class in Python 2.7...
300	distutils.command.clean.clean.run(self)	300	distutils.command.clean.clean.run(self)
301		301
302	remove_files = [	302	remove_files = [
303	'torch_npu/csrc/aten/RegisterCPU.cpp',	303	'torch_npu/csrc/aten/RegisterCPU.cpp',
304	'torch_npu/csrc/aten/RegisterNPU.cpp',	304	'torch_npu/csrc/aten/RegisterNPU.cpp',
305	'torch_npu/csrc/aten/RegisterAutogradNPU.cpp',	305	'torch_npu/csrc/aten/RegisterAutogradNPU.cpp',
306	'torch_npu/csrc/aten/NPUNativeFunctions.h',	306	'torch_npu/csrc/aten/NPUNativeFunctions.h',
307	'torch_npu/csrc/aten/CustomRegisterSchema.cpp',	307	'torch_npu/csrc/aten/CustomRegisterSchema.cpp',
308	'torch_npu/csrc/aten/ForeachRegister.cpp',	308	'torch_npu/csrc/aten/ForeachRegister.cpp',
309	'torch_npu/utils/custom_ops.py',	309	'torch_npu/utils/custom_ops.py',
310	'torch_npu/version.py',	310	'torch_npu/version.py',
311	]	311	]
312	for remove_file in remove_files:	312	for remove_file in remove_files:
313	file_path = os.path.join(BASE_DIR, remove_file)	313	file_path = os.path.join(BASE_DIR, remove_file)
314	if os.path.exists(file_path):	314	if os.path.exists(file_path):
315	os.remove(file_path)	315	os.remove(file_path)
316		316
317	USE_NINJA = os.environ["CMAKE_GENERATOR"].lower() == "ninja" if "CMAKE_GENERATOR" in os.environ else shutil.which("ninja")	317	USE_NINJA = os.environ["CMAKE_GENERATOR"].lower() == "ninja" if "CMAKE_GENERATOR" in os.environ else shutil.which("ninja")
318		318
319	class CPPLibBuild(build_clib, object):	319	class CPPLibBuild(build_clib, object):
320	def run(self):	320	def run(self):
321	cmake = get_cmake_command()	321	cmake = get_cmake_command()
322		322
323	if cmake is None:	323	if cmake is None:
324	raise RuntimeError(	324	raise RuntimeError(
325	"CMake must be installed to build the following extensions: " +	325	"CMake must be installed to build the following extensions: " +
326	", ".join(e.name for e in self.extensions))	326	", ".join(e.name for e in self.extensions))
327	self.cmake = cmake	327	self.cmake = cmake
328		328
329	build_dir = os.path.join(BASE_DIR, "build")	329	build_dir = os.path.join(BASE_DIR, "build")
330	build_type_dir = os.path.join(build_dir)	330	build_type_dir = os.path.join(build_dir)
331	output_lib_path = os.path.join(build_type_dir, "packages/torch_npu/lib")	331	output_lib_path = os.path.join(build_type_dir, "packages/torch_npu/lib")
332	os.makedirs(build_type_dir, exist_ok=True)	332	os.makedirs(build_type_dir, exist_ok=True)
333	os.chmod(build_type_dir, mode=BUILD_PERMISSION)	333	os.chmod(build_type_dir, mode=BUILD_PERMISSION)
334	os.makedirs(output_lib_path, exist_ok=True)	334	os.makedirs(output_lib_path, exist_ok=True)
335	self.build_lib = os.path.relpath(os.path.join(build_dir, "packages/torch_npu"))	335	self.build_lib = os.path.relpath(os.path.join(build_dir, "packages/torch_npu"))
336	self.build_temp = os.path.relpath(build_type_dir)	336	self.build_temp = os.path.relpath(build_type_dir)
337		337
338	cmake_args = [	338	cmake_args = [
339	'-DCMAKE_BUILD_TYPE=' + get_build_type(),	339	'-DCMAKE_BUILD_TYPE=' + get_build_type(),
340	'-DCMAKE_INSTALL_PREFIX=' + os.path.realpath(output_lib_path),	340	'-DCMAKE_INSTALL_PREFIX=' + os.path.realpath(output_lib_path),
341	'-DCMAKE_LIBRARY_OUTPUT_DIRECTORY=' + os.path.realpath(output_lib_path),	341	'-DCMAKE_LIBRARY_OUTPUT_DIRECTORY=' + os.path.realpath(output_lib_path),
342	'-DCMAKE_ARCHIVE_OUTPUT_DIRECTORY=' + os.path.realpath(output_lib_path),	342	'-DCMAKE_ARCHIVE_OUTPUT_DIRECTORY=' + os.path.realpath(output_lib_path),
343	'-DTORCHNPU_INSTALL_LIBDIR=' + os.path.realpath(output_lib_path),	343	'-DTORCHNPU_INSTALL_LIBDIR=' + os.path.realpath(output_lib_path),
344	'-DPYTHON_INCLUDE_DIR=' + get_paths().get('include'),	344	'-DPYTHON_INCLUDE_DIR=' + get_paths().get('include'),
345	'-DTORCH_VERSION=' + VERSION,	345	'-DTORCH_VERSION=' + VERSION,
346	'-DPYTORCH_INSTALL_DIR=' + get_pytorch_dir()]	346	'-DPYTORCH_INSTALL_DIR=' + get_pytorch_dir()]
347		347
348	if DISABLE_TORCHAIR == 'FALSE':	348	if DISABLE_TORCHAIR == 'FALSE':
349	if check_torchair_valid(BASE_DIR):	349	if check_torchair_valid(BASE_DIR):
350	cmake_args.append('-DBUILD_TORCHAIR=on')	350	cmake_args.append('-DBUILD_TORCHAIR=on')
351	torchair_install_prefix = os.path.join(build_type_dir, "packages/torch_npu/dynamo/torchair")	351	torchair_install_prefix = os.path.join(build_type_dir, "packages/torch_npu/dynamo/torchair")
352	cmake_args.append(f'-DTORCHAIR_INSTALL_PREFIX={torchair_install_prefix}')	352	cmake_args.append(f'-DTORCHAIR_INSTALL_PREFIX={torchair_install_prefix}')
353	cmake_args.append(f'-DTORCHAIR_TARGET_PYTHON={sys.executable}')	353	cmake_args.append(f'-DTORCHAIR_TARGET_PYTHON={sys.executable}')
354		354
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
367	if os.getenv('_ABI_VERSION') is not None:	367	if os.getenv('_ABI_VERSION') is not None:
368	cmake_args.append('-DABI_VERSION=' + os.getenv('_ABI_VERSION'))	368	cmake_args.append('-DABI_VERSION=' + os.getenv('_ABI_VERSION'))
369		369
370	if USE_NINJA:	370	if USE_NINJA:
371	cmake_args.append("-GNinja")	371	cmake_args.append("-GNinja")
372		372
373	max_jobs = os.getenv("MAX_JOBS")	373	max_jobs = os.getenv("MAX_JOBS")
374	if max_jobs is not None or not USE_NINJA:	374	if max_jobs is not None or not USE_NINJA:
375	max_jobs = max_jobs or str(multiprocessing.cpu_count())	375	max_jobs = max_jobs or str(multiprocessing.cpu_count())
376	build_args = ['-j', max_jobs]	376	build_args = ['-j', max_jobs]
377	else:	377	else:
378	build_args = []	378	build_args = []
379		379
380	subprocess.check_call([self.cmake, BASE_DIR] + cmake_args, cwd=build_type_dir, env=os.environ)	380	subprocess.check_call([self.cmake, BASE_DIR] + cmake_args, cwd=build_type_dir, env=os.environ)
381	for base_dir, dirs, files in os.walk(build_type_dir):	381	for base_dir, dirs, files in os.walk(build_type_dir):
382	for dir_name in dirs:	382	for dir_name in dirs:
383	dir_path = os.path.join(base_dir, dir_name)	383	dir_path = os.path.join(base_dir, dir_name)
384	os.chmod(dir_path, mode=BUILD_PERMISSION)	384	os.chmod(dir_path, mode=BUILD_PERMISSION)
385	for file_name in files:	385	for file_name in files:
386	file_path = os.path.join(base_dir, file_name)	386	file_path = os.path.join(base_dir, file_name)
387	os.chmod(file_path, mode=BUILD_PERMISSION)	387	os.chmod(file_path, mode=BUILD_PERMISSION)
388		388
389	if USE_NINJA:	389	if USE_NINJA:
390	subprocess.check_call(['ninja'] + build_args, cwd=build_type_dir, env=os.environ)	390	subprocess.check_call(['ninja'] + build_args, cwd=build_type_dir, env=os.environ)
391	else:	391	else:
392	subprocess.check_call(['make'] + build_args, cwd=build_type_dir, env=os.environ)	392	subprocess.check_call(['make'] + build_args, cwd=build_type_dir, env=os.environ)
393		393
394		394
395	class Build(build_ext, object):	395	class Build(build_ext, object):
396		396
397	def run(self):	397	def run(self):
398	self.run_command('build_clib')	398	self.run_command('build_clib')
399	self.run_command('build_py')	399	self.run_command('build_py')
400	# proceed with the normal build_ext process	400	# proceed with the normal build_ext process
401	self.build_lib = os.path.relpath(os.path.join(BASE_DIR, "build/packages"))	401	self.build_lib = os.path.relpath(os.path.join(BASE_DIR, "build/packages"))
402	self.build_temp = os.path.relpath(os.path.join(BASE_DIR, "build/temp"))	402	self.build_temp = os.path.relpath(os.path.join(BASE_DIR, "build/temp"))
403	self.library_dirs.append(	403	self.library_dirs.append(
404	os.path.relpath(os.path.join(BASE_DIR, "build/packages/torch_npu/lib"))	404	os.path.relpath(os.path.join(BASE_DIR, "build/packages/torch_npu/lib"))
405	)	405	)
406	super(Build, self).run()	406	super(Build, self).run()
407		407
408		408
409	class InstallCmd(install):	409	class InstallCmd(install):
410		410
411	def finalize_options(self) -> None:	411	def finalize_options(self) -> None:
412	self.build_lib = os.path.relpath(os.path.join(BASE_DIR, "build/packages"))	412	self.build_lib = os.path.relpath(os.path.join(BASE_DIR, "build/packages"))
413	return super(InstallCmd, self).finalize_options()	413	return super(InstallCmd, self).finalize_options()
414		414
415		415
416	def add_ops_files(base_dir, file_list):	416	def add_ops_files(base_dir, file_list):
417	# add ops header files	417	# add ops header files
418	plugin_path = os.path.join(base_dir, 'third_party/op-plugin/op_plugin/include')	418	plugin_path = os.path.join(base_dir, 'third_party/op-plugin/op_plugin/include')
419	if os.path.exists(plugin_path):	419	if os.path.exists(plugin_path):
420	file_list.append('third_party/op-plugin/op_plugin/include/*.h')	420	file_list.append('third_party/op-plugin/op_plugin/include/*.h')
421	plugin_utils_path = os.path.join(base_dir, 'third_party/op-plugin/op_plugin/utils')	421	plugin_utils_path = os.path.join(base_dir, 'third_party/op-plugin/op_plugin/utils')
422	if os.path.exists(plugin_utils_path):	422	if os.path.exists(plugin_utils_path):
423	file_list.append('third_party/op-plugin/op_plugin/utils/*.h')	423	file_list.append('third_party/op-plugin/op_plugin/utils/*.h')
424	return	424	return
425		425
426		426
427	def add_ops_python_files(ret_list):	427	def add_ops_python_files(ret_list):
428	# add ops python files	428	# add ops python files
429	opplugin_path = os.path.join(BASE_DIR, 'third_party/op-plugin/op_plugin/python')	429	opplugin_path = os.path.join(BASE_DIR, 'third_party/op-plugin/op_plugin/python')
430		430
431	if os.path.exists(opplugin_path):	431	if os.path.exists(opplugin_path):
432	ops_python_files = glob.glob(os.path.join(opplugin_path, '*/.py'), recursive=True)	432	ops_python_files = glob.glob(os.path.join(opplugin_path, '*/.py'), recursive=True)
433	for src in ops_python_files:	433	for src in ops_python_files:
434	dst = os.path.join(	434	dst = os.path.join(
435	os.path.join(BASE_DIR, "build/packages/torch_npu/op_plugin"),	435	os.path.join(BASE_DIR, "build/packages/torch_npu/op_plugin"),
436	os.path.relpath(src, opplugin_path))	436	os.path.relpath(src, opplugin_path))
437	os.makedirs(os.path.dirname(dst), exist_ok=True)	437	os.makedirs(os.path.dirname(dst), exist_ok=True)
438	ret_list.append((src, dst))	438	ret_list.append((src, dst))
439	return	439	return
440		440
441		441
442	def get_src_py_and_dst():	442	def get_src_py_and_dst():
443	ret = []	443	ret = []
444	generated_python_files = glob.glob(	444	generated_python_files = glob.glob(
445	os.path.join(BASE_DIR, "torch_npu", '*/.py'),	445	os.path.join(BASE_DIR, "torch_npu", '*/.py'),
446	recursive=True) + glob.glob(	446	recursive=True) + glob.glob(
447	os.path.join(BASE_DIR, "torch_npu", '*/.yaml'),	447	os.path.join(BASE_DIR, "torch_npu", '*/.yaml'),
448	recursive=True) + glob.glob(	448	recursive=True) + glob.glob(
449	os.path.join(BASE_DIR, "torch_npu", 'acl*.json'),	449	os.path.join(BASE_DIR, "torch_npu", 'acl*.json'),
450	recursive=True) + glob.glob(	450	recursive=True) + glob.glob(
451	os.path.join(BASE_DIR, "torch_npu", 'contrib/apis_config.json'),	451	os.path.join(BASE_DIR, "torch_npu", 'contrib/apis_config.json'),
452	recursive=True)	452	recursive=True)
453	for src in generated_python_files:	453	for src in generated_python_files:
454	dst = os.path.join(	454	dst = os.path.join(
455	os.path.join(BASE_DIR, "build/packages/torch_npu"),	455	os.path.join(BASE_DIR, "build/packages/torch_npu"),
456	os.path.relpath(src, os.path.join(BASE_DIR, "torch_npu")))	456	os.path.relpath(src, os.path.join(BASE_DIR, "torch_npu")))
457	os.makedirs(os.path.dirname(dst), exist_ok=True)	457	os.makedirs(os.path.dirname(dst), exist_ok=True)
458	ret.append((src, dst))	458	ret.append((src, dst))
459		459
460	add_ops_python_files(ret)	460	add_ops_python_files(ret)
461		461
462	header_files = [	462	header_files = [
463	"torch_npu/csrc/*.h",	463	"torch_npu/csrc/*.h",
464	"torch_npu/csrc//.h",	464	"torch_npu/csrc//.h",
465	"torch_npu/csrc//.hpp",	465	"torch_npu/csrc//.hpp",
466	"torch_npu/csrc///*.h",	466	"torch_npu/csrc///*.h",
467	"torch_npu/csrc////.h",	467	"torch_npu/csrc////.h",
468	"torch_npu/csrc/////*.h",	468	"torch_npu/csrc/////*.h",
469	"third_party/acl/inc//.h",	469	"third_party/acl/inc//.h",
470	"third_party/hccl/inc//.h",	470	"third_party/hccl/inc//.h",
471	"third_party/acl/inc///*.h",	471	"third_party/acl/inc///*.h",
472	"torch_npu/csrc/distributed/HCCLUtils.hpp",	472	"torch_npu/csrc/distributed/HCCLUtils.hpp",
473	"torch_npu/csrc/distributed/ProcessGroupHCCL.hpp"	473	"torch_npu/csrc/distributed/ProcessGroupHCCL.hpp"
474	]	474	]
475	add_ops_files(BASE_DIR, header_files)	475	add_ops_files(BASE_DIR, header_files)
476	glob_header_files = []	476	glob_header_files = []
477	for regex_pattern in header_files:	477	for regex_pattern in header_files:
478	glob_header_files += glob.glob(os.path.join(BASE_DIR, regex_pattern), recursive=True)	478	glob_header_files += glob.glob(os.path.join(BASE_DIR, regex_pattern), recursive=True)
479		479
480	for src in glob_header_files:	480	for src in glob_header_files:
481	dst = os.path.join(	481	dst = os.path.join(
482	os.path.join(BASE_DIR, "build/packages/torch_npu/include/torch_npu"),	482	os.path.join(BASE_DIR, "build/packages/torch_npu/include/torch_npu"),
483	os.path.relpath(src, os.path.join(BASE_DIR, "torch_npu")))	483	os.path.relpath(src, os.path.join(BASE_DIR, "torch_npu")))
484	os.makedirs(os.path.dirname(dst), exist_ok=True)	484	os.makedirs(os.path.dirname(dst), exist_ok=True)
485	ret.append((src, dst))	485	ret.append((src, dst))
486		486
487	torch_header_files = [	487	torch_header_files = [
488	"/.h",	488	"/.h",
489	"//*.h",	489	"//*.h",
490	"///.h",	490	"///.h",
491	"////*.h",	491	"////*.h",
492	"/////.h"	492	"/////.h"
493	]	493	]
494	torch_glob_header_files = []	494	torch_glob_header_files = []
495	for regex_pattern in torch_header_files:	495	for regex_pattern in torch_header_files:
496	torch_glob_header_files += glob.glob(os.path.join(BASE_DIR, "patch/include", regex_pattern), recursive=True)	496	torch_glob_header_files += glob.glob(os.path.join(BASE_DIR, "patch/include", regex_pattern), recursive=True)
497		497
498	for src in torch_glob_header_files:	498	for src in torch_glob_header_files:
499	dst = os.path.join(	499	dst = os.path.join(
500	os.path.join(BASE_DIR, "build/packages/torch_npu/include"),	500	os.path.join(BASE_DIR, "build/packages/torch_npu/include"),
501	os.path.relpath(src, os.path.join(BASE_DIR, "patch/include")))	501	os.path.relpath(src, os.path.join(BASE_DIR, "patch/include")))
502	os.makedirs(os.path.dirname(dst), exist_ok=True)	502	os.makedirs(os.path.dirname(dst), exist_ok=True)
503	ret.append((src, dst))	503	ret.append((src, dst))
504		504
505	aot_inductor_files = [	505	aot_inductor_files = [
506	# Follow torch v2.6.0.	506	# Follow torch v2.6.0.
507	# These aoti_runtime/*.cpp don't compile to libtorch_npu,	507	# These aoti_runtime/*.cpp don't compile to libtorch_npu,
508	# but act like header files when generate cppwrapper in aot-inductor.	508	# but act like header files when generate cppwrapper in aot-inductor.
509	"torch_npu/_inductor/codegen/aoti_runtime/*.cpp"	509	"torch_npu/_inductor/codegen/aoti_runtime/*.cpp"
510	]	510	]
511	glob_aoti_files = []	511	glob_aoti_files = []
512	for regex_pattern in aot_inductor_files:	512	for regex_pattern in aot_inductor_files:
513	glob_aoti_files += glob.glob(	513	glob_aoti_files += glob.glob(
514	os.path.join(BASE_DIR, regex_pattern), recursive=True	514	os.path.join(BASE_DIR, regex_pattern), recursive=True
515	)	515	)
516		516
517	for src in glob_aoti_files:	517	for src in glob_aoti_files:
518	# Dst: torch_npu/_inductor/codegen/aoti_runtime/*.cpp	518	# Dst: torch_npu/_inductor/codegen/aoti_runtime/*.cpp
519	dst = os.path.join(	519	dst = os.path.join(
520	os.path.join(BASE_DIR, "build/packages/torch_npu/"),	520	os.path.join(BASE_DIR, "build/packages/torch_npu/"),
521	os.path.relpath(src, os.path.join(BASE_DIR, "torch_npu")),	521	os.path.relpath(src, os.path.join(BASE_DIR, "torch_npu")),
522	)	522	)
523	os.makedirs(os.path.dirname(dst), exist_ok=True)	523	os.makedirs(os.path.dirname(dst), exist_ok=True)
524	ret.append((src, dst))	524	ret.append((src, dst))
525		525
526	def add_torch_npu_codegen(codegen_src_dir, codegen_dst_dir, exclude_root_init=None):	526	def add_torch_npu_codegen(codegen_src_dir, codegen_dst_dir, exclude_root_init=None):
527	"""	527	"""
528	复制codegen目录下的文件到目标路径	528	复制codegen目录下的文件到目标路径
529	:param codegen_src_dir: 源codegen目录	529	:param codegen_src_dir: 源codegen目录
530	:param codegen_dst_dir: 目标目录	530	:param codegen_dst_dir: 目标目录
531	:param exclude_root_init: 需要排除根目录__init__.py的源目录(仅过滤该目录下的__init__.py)	531	:param exclude_root_init: 需要排除根目录__init__.py的源目录(仅过滤该目录下的__init__.py)
532	"""	532	"""
533	# 匹配需要复制的文件类型	533	# 匹配需要复制的文件类型
534	codegen_files = glob.glob(	534	codegen_files = glob.glob(
535	os.path.join(codegen_src_dir, '*/.py'), recursive=True	535	os.path.join(codegen_src_dir, '*/.py'), recursive=True
536	) + glob.glob(	536	) + glob.glob(
537	os.path.join(codegen_src_dir, '*/.yaml'), recursive=True	537	os.path.join(codegen_src_dir, '*/.yaml'), recursive=True
538	) + glob.glob(	538	) + glob.glob(
539	os.path.join(codegen_src_dir, '*/.json'), recursive=True	539	os.path.join(codegen_src_dir, '*/.json'), recursive=True
540	) + glob.glob(	540	) + glob.glob(
541	os.path.join(codegen_src_dir, '*/.cpp'), recursive=True	541	os.path.join(codegen_src_dir, '*/.cpp'), recursive=True
542	) + glob.glob(	542	) + glob.glob(
543	os.path.join(codegen_src_dir, '*/.h'), recursive=True	543	os.path.join(codegen_src_dir, '*/.h'), recursive=True
544	)	544	)
545		545
546	# 按原目录结构复制到目标路径	546	# 按原目录结构复制到目标路径
547	for src in codegen_files:	547	for src in codegen_files:
548	# 仅过滤指定目录下的根级__init__.py	548	# 仅过滤指定目录下的根级__init__.py
549	if (exclude_root_init is not None and	549	if (exclude_root_init is not None and
550	os.path.basename(src) == '__init__.py' and	550	os.path.basename(src) == '__init__.py' and
551	os.path.dirname(src) == exclude_root_init):	551	os.path.dirname(src) == exclude_root_init):
552	continue # 跳过op-plugin/codegen根目录的__init__.py	552	continue # 跳过op-plugin/codegen根目录的__init__.py
553		553
554	# 计算目标路径（保留原目录层级）	554	# 计算目标路径（保留原目录层级）
555	dst = os.path.join(	555	dst = os.path.join(
556	codegen_dst_dir,	556	codegen_dst_dir,
557	os.path.relpath(src, codegen_src_dir) # 保留torchnpugen内部的目录层级	557	os.path.relpath(src, codegen_src_dir) # 保留torchnpugen内部的目录层级
558	)	558	)
559	print(os.path.relpath(src, codegen_src_dir))	559	print(os.path.relpath(src, codegen_src_dir))
560	# 确保目标目录存在	560	# 确保目标目录存在
561	os.makedirs(os.path.dirname(dst), exist_ok=True)	561	os.makedirs(os.path.dirname(dst), exist_ok=True)
562	# 加入文件复制列表	562	# 加入文件复制列表
563	ret.append((src, dst))	563	ret.append((src, dst))
564		564
565	# 新增：提前创建 torchnpugen 根目录	565	# 新增：提前创建 torchnpugen 根目录
566	torchnpugen_root = os.path.join(BASE_DIR, "build/packages/torchnpugen")	566	torchnpugen_root = os.path.join(BASE_DIR, "build/packages/torchnpugen")
567	os.makedirs(torchnpugen_root, exist_ok=True)	567	os.makedirs(torchnpugen_root, exist_ok=True)
568	# 将codegen复制到package路径	568	# 将codegen复制到package路径
569	codegen_src_dir = os.path.join(BASE_DIR, "torchnpugen")	569	codegen_src_dir = os.path.join(BASE_DIR, "torchnpugen")
570	codegen_dst_dir = os.path.join(BASE_DIR, "build/packages/torchnpugen")	570	codegen_dst_dir = os.path.join(BASE_DIR, "build/packages/torchnpugen")
571	# 复制torch_npu的torchnpugen	571	# 复制torch_npu的torchnpugen
572	add_torch_npu_codegen(codegen_src_dir, codegen_dst_dir)	572	add_torch_npu_codegen(codegen_src_dir, codegen_dst_dir)
573	# 复制op-plugin的torchnpugen（仅过滤其根目录的__init__.py）	573	# 复制op-plugin的torchnpugen（仅过滤其根目录的__init__.py）
574	op_plugin_codegen_src = os.path.join(BASE_DIR, "third_party/op-plugin/torchnpugen")	574	op_plugin_codegen_src = os.path.join(BASE_DIR, "third_party/op-plugin/torchnpugen")
575	add_torch_npu_codegen(	575	add_torch_npu_codegen(
576	op_plugin_codegen_src,	576	op_plugin_codegen_src,
577	codegen_dst_dir,	577	codegen_dst_dir,
578	exclude_root_init=op_plugin_codegen_src # 指定要过滤根目录__init__.py的源目录	578	exclude_root_init=op_plugin_codegen_src # 指定要过滤根目录__init__.py的源目录
579	)	579	)
580		580
581	return ret	581	return ret
582		582
583		583
584	class EggInfoBuild(egg_info, object):	584	class EggInfoBuild(egg_info, object):
585	def finalize_options(self):	585	def finalize_options(self):
586	self.egg_base = os.path.relpath(os.path.join(BASE_DIR, "build/packages"))	586	self.egg_base = os.path.relpath(os.path.join(BASE_DIR, "build/packages"))
587	ret = get_src_py_and_dst()	587	ret = get_src_py_and_dst()
588	for src, dst in ret:	588	for src, dst in ret:
589	self.copy_file(src, dst)	589	self.copy_file(src, dst)
590	super(EggInfoBuild, self).finalize_options()	590	super(EggInfoBuild, self).finalize_options()
591		591
592		592
593	class PythonPackageBuild(build_py, object):	593	class PythonPackageBuild(build_py, object):
594	def run(self) -> None:	594	def run(self) -> None:
595	ret = get_src_py_and_dst()	595	ret = get_src_py_and_dst()
596	for src, dst in ret:	596	for src, dst in ret:
597	self.copy_file(src, dst)	597	self.copy_file(src, dst)
598	super(PythonPackageBuild, self).finalize_options()	598	super(PythonPackageBuild, self).finalize_options()
599		599
600		600
601	class BdistWheelBuild(bdist_wheel):	601	class BdistWheelBuild(bdist_wheel):
602	def run(self):	602	def run(self):
603	if which('patchelf') is not None:	603	if which('patchelf') is not None:
604	patchelf_dynamic_library()	604	patchelf_dynamic_library()
605		605
606	if not DEBUG and which('eu-strip') is not None:	606	if not DEBUG and which('eu-strip') is not None:
607	generate_dbg_files_and_strip()	607	generate_dbg_files_and_strip()
608		608
609	torch_dependencies = ["libc10.so", "libtorch.so", "libtorch_cpu.so", "libtorch_python.so"]	609	torch_dependencies = ["libc10.so", "libtorch.so", "libtorch_cpu.so", "libtorch_python.so"]
610	cann_dependencies = ["libhccl.so", "libascendcl.so", "libacl_op_compiler.so", "libge_runner.so",	610	cann_dependencies = ["libhccl.so", "libascendcl.so", "libacl_op_compiler.so", "libge_runner.so",
611	"libgraph.so", "libacl_tdt_channel.so", "libfmk_parser.so", "libascend_protobuf.so",	611	"libgraph.so", "libacl_tdt_channel.so", "libfmk_parser.so", "libascend_protobuf.so",
612	"libascend_ml.so"]	612	"libascend_ml.so"]
613	other_dependencies = ["libtorch_npu.so", "libnpu_profiler.so", "libgomp.so.1", "libatb.so"]	613	other_dependencies = ["libtorch_npu.so", "libnpu_profiler.so", "libgomp.so.1", "libatb.so"]
614		614
615	dependencies = torch_dependencies + cann_dependencies + other_dependencies	615	dependencies = torch_dependencies + cann_dependencies + other_dependencies
616		616
617	bdist_wheel.run(self)	617	bdist_wheel.run(self)
618		618
619	if is_manylinux:	619	if is_manylinux:
620	file = glob.glob(os.path.join(self.dist_dir, "linux.whl"))[0]	620	file = glob.glob(os.path.join(self.dist_dir, "linux.whl"))[0]
621		621
622	auditwheel_cmd = ["auditwheel", "-v", "repair", "-w", self.dist_dir, file]	622	auditwheel_cmd = ["auditwheel", "-v", "repair", "-w", self.dist_dir, file]
623	for i in dependencies:	623	for i in dependencies:
624	auditwheel_cmd += ["--exclude", i]	624	auditwheel_cmd += ["--exclude", i]
625		625
626	try:	626	try:
627	subprocess.run(auditwheel_cmd, check=True, stdout=subprocess.PIPE)	627	subprocess.run(auditwheel_cmd, check=True, stdout=subprocess.PIPE)
628	finally:	628	finally:
629	os.remove(file)	629	os.remove(file)
630		630
631		631
632	build_mode = _get_build_mode()	632	build_mode = _get_build_mode()
633	if build_mode not in ['clean']:	633	if build_mode not in ['clean']:
634	# Generate bindings code, including RegisterNPU.cpp & NPUNativeFunctions.h.	634	# Generate bindings code, including RegisterNPU.cpp & NPUNativeFunctions.h.
635	generate_bindings_code(BASE_DIR)	635	generate_bindings_code(BASE_DIR)
636	if Path(BASE_DIR).joinpath("third_party/Tensorpipe/third_party/acl/libs").exists():	636	if Path(BASE_DIR).joinpath("third_party/Tensorpipe/third_party/acl/libs").exists():
637	build_stub(Path(BASE_DIR).joinpath("third_party/Tensorpipe"))	637	build_stub(Path(BASE_DIR).joinpath("third_party/Tensorpipe"))
638	build_stub(BASE_DIR)	638	build_stub(BASE_DIR)
639		639
640	# Setup include directories folders.	640	# Setup include directories folders.
641	include_directories = [	641	include_directories = [
642	BASE_DIR,	642	BASE_DIR,
643	os.path.join(BASE_DIR, 'patch/include'),	643	os.path.join(BASE_DIR, 'patch/include'),
644	os.path.join(BASE_DIR, 'third_party/hccl/inc'),	644	os.path.join(BASE_DIR, 'third_party/hccl/inc'),
645	os.path.join(BASE_DIR, 'third_party/acl/inc')	645	os.path.join(BASE_DIR, 'third_party/acl/inc')
646	]	646	]
647		647
648	extra_link_args = []	648	extra_link_args = []
649		649
650	DEBUG = (os.getenv('DEBUG', default='').upper() in ['ON', '1', 'YES', 'TRUE', 'Y'])	650	DEBUG = (os.getenv('DEBUG', default='').upper() in ['ON', '1', 'YES', 'TRUE', 'Y'])
651		651
652	extra_compile_args = [	652	extra_compile_args = [
653	'-std=c++17',	653	'-std=c++17',
654	'-Wno-sign-compare',	654	'-Wno-sign-compare',
655	'-Wno-deprecated-declarations',	655	'-Wno-deprecated-declarations',
656	'-Wno-return-type'	656	'-Wno-return-type'
657	]	657	]
658		658
659	if re.match(r'clang', os.getenv('CC', '')):	659	if re.match(r'clang', os.getenv('CC', '')):
660	extra_compile_args += [	660	extra_compile_args += [
661	'-Wno-macro-redefined',	661	'-Wno-macro-redefined',
662	'-Wno-return-std-move',	662	'-Wno-return-std-move',
663	]	663	]
664		664
665	if DEBUG:	665	if DEBUG:
666	extra_compile_args += ['-O0', '-g']	666	extra_compile_args += ['-O0', '-g']
667	extra_link_args += ['-O0', '-g', '-Wl,-z,now']	667	extra_link_args += ['-O0', '-g', '-Wl,-z,now']
668	else:	668	else:
669	extra_compile_args += ['-DNDEBUG']	669	extra_compile_args += ['-DNDEBUG']
670	extra_link_args += ['-Wl,-z,now']	670	extra_link_args += ['-Wl,-z,now']
671		671
672	# valid manylinux tags	672	# valid manylinux tags
673	manylinux_tags = [	673	manylinux_tags = [
674	"manylinux1_x86_64",	674	"manylinux1_x86_64",
675	"manylinux2010_x86_64",	675	"manylinux2010_x86_64",
676	"manylinux2014_x86_64",	676	"manylinux2014_x86_64",
677	"manylinux2014_aarch64",	677	"manylinux2014_aarch64",
678	"manylinux_2_5_x86_64",	678	"manylinux_2_5_x86_64",
679	"manylinux_2_12_x86_64",	679	"manylinux_2_12_x86_64",
680	"manylinux_2_17_x86_64",	680	"manylinux_2_17_x86_64",
681	"manylinux_2_17_aarch64",	681	"manylinux_2_17_aarch64",
682	"manylinux_2_24_x86_64",	682	"manylinux_2_24_x86_64",
683	"manylinux_2_24_aarch64",	683	"manylinux_2_24_aarch64",
684	"manylinux_2_27_x86_64",	684	"manylinux_2_27_x86_64",
685	"manylinux_2_27_aarch64",	685	"manylinux_2_27_aarch64",
686	"manylinux_2_28_x86_64",	686	"manylinux_2_28_x86_64",
687	"manylinux_2_28_aarch64",	687	"manylinux_2_28_aarch64",
688	"manylinux_2_31_x86_64",	688	"manylinux_2_31_x86_64",
689	"manylinux_2_31_aarch64",	689	"manylinux_2_31_aarch64",
690	"manylinux_2_34_x86_64",	690	"manylinux_2_34_x86_64",
691	"manylinux_2_34_aarch64",	691	"manylinux_2_34_aarch64",
692	"manylinux_2_35_x86_64"	692	"manylinux_2_35_x86_64"
693	"manylinux_2_35_aarch64",	693	"manylinux_2_35_aarch64",
694	]	694	]
695	is_manylinux = os.environ.get("AUDITWHEEL_PLAT", None) in manylinux_tags	695	is_manylinux = os.environ.get("AUDITWHEEL_PLAT", None) in manylinux_tags
696		696
697	readme = os.path.join(BASE_DIR, "README.md")	697	readme = os.path.join(BASE_DIR, "README.md")
698	if not os.path.exists(readme):	698	if not os.path.exists(readme):
699	raise FileNotFoundError("Unable to find 'README.md'")	699	raise FileNotFoundError("Unable to find 'README.md'")
700	with open(readme, encoding="utf-8") as fdesc:	700	with open(readme, encoding="utf-8") as fdesc:
701	long_description = fdesc.read()	701	long_description = fdesc.read()
702		702
703	classifiers = [	703	classifiers = [
704	"Development Status :: 5 - Production/Stable",	704	"Development Status :: 5 - Production/Stable",
705	"Intended Audience :: Developers",	705	"Intended Audience :: Developers",
706	"License :: OSI Approved :: BSD License",	706	"License :: OSI Approved :: BSD License",
707	"Operating System :: POSIX :: Linux",	707	"Operating System :: POSIX :: Linux",
708	"Topic :: Scientific/Engineering",	708	"Topic :: Scientific/Engineering",
709	"Topic :: Scientific/Engineering :: Mathematics",	709	"Topic :: Scientific/Engineering :: Mathematics",
710	"Topic :: Scientific/Engineering :: Artificial Intelligence",	710	"Topic :: Scientific/Engineering :: Artificial Intelligence",
711	"Topic :: Software Development",	711	"Topic :: Software Development",
712	"Topic :: Software Development :: Libraries",	712	"Topic :: Software Development :: Libraries",
713	"Topic :: Software Development :: Libraries :: Python Modules",	713	"Topic :: Software Development :: Libraries :: Python Modules",
714	"Programming Language :: Python",	714	"Programming Language :: Python",
715	"Programming Language :: Python :: 3 :: Only",	715	"Programming Language :: Python :: 3 :: Only",
716	"Programming Language :: Python :: 3.8",	716	"Programming Language :: Python :: 3.8",
717	"Programming Language :: Python :: 3.9",	717	"Programming Language :: Python :: 3.9",
718	"Programming Language :: Python :: 3.10",	718	"Programming Language :: Python :: 3.10",
719	"Programming Language :: Python :: 3.11",	719	"Programming Language :: Python :: 3.11",
720	]	720	]
721		721
722	requirements = ['torch==2.7.1+cpu' if platform.machine() == 'x86_64' else 'torch==2.7.1']	722	requirements = ['torch==2.7.1+cpu' if platform.machine() == 'x86_64' else 'torch==2.7.1']
723		723
724	ext_modules = [CppExtension(	724	ext_modules = [CppExtension(
725	'torch_npu._C',	725	'torch_npu._C',
726	sources=["torch_npu/csrc/InitNpuBindings.cpp"],	726	sources=["torch_npu/csrc/InitNpuBindings.cpp"],
727	libraries=["torch_npu"],	727	libraries=["torch_npu"],
728	include_dirs=include_directories,	728	include_dirs=include_directories,
729	extra_compile_args=extra_compile_args + ['-fstack-protector-all'] + [	729	extra_compile_args=extra_compile_args + ['-fstack-protector-all'] + [
730	'-D__FILENAME__=\"InitNpuBindings.cpp\"'],	730	'-D__FILENAME__=\"InitNpuBindings.cpp\"'],
731	library_dirs=["lib"],	731	library_dirs=["lib"],
732	extra_link_args=extra_link_args + ['-Wl,-rpath,$ORIGIN/lib', '-Wl,-Bsymbolic-functions'],	732	extra_link_args=extra_link_args + ['-Wl,-rpath,$ORIGIN/lib', '-Wl,-Bsymbolic-functions'],
733	define_macros=[('_GLIBCXX_USE_CXX11_ABI', '1' if USE_CXX11_ABI else '0'),	733	define_macros=[('_GLIBCXX_USE_CXX11_ABI', '1' if USE_CXX11_ABI else '0'),
734	('GLIBCXX_USE_CXX11_ABI', '1' if USE_CXX11_ABI else '0')]	734	('GLIBCXX_USE_CXX11_ABI', '1' if USE_CXX11_ABI else '0')]
735	)]	735	)]
736		736
737	setup(	737	setup(
738	name=os.environ.get('TORCH_NPU_PACKAGE_NAME', 'torch_npu'),	738	name=os.environ.get('TORCH_NPU_PACKAGE_NAME', 'torch_npu'),
739	version=VERSION,	739	version=VERSION,
740	description='NPU bridge for PyTorch',	740	description='NPU bridge for PyTorch',
741	long_description=long_description,	741	long_description=long_description,
742	long_description_content_type="text/markdown",	742	long_description_content_type="text/markdown",
743	license="BSD License",	743	license="BSD License",
744	classifiers=classifiers,	744	classifiers=classifiers,
745	packages=["torch_npu", "torchnpugen"],	745	packages=["torch_npu", "torchnpugen"],
746	libraries=[('torch_npu', {'sources': list()})],	746	libraries=[('torch_npu', {'sources': list()})],
747	package_dir={'': os.path.relpath(os.path.join(BASE_DIR, "build/packages"))},	747	package_dir={'': os.path.relpath(os.path.join(BASE_DIR, "build/packages"))},
748	ext_modules=ext_modules,	748	ext_modules=ext_modules,
749	install_requires=requirements,	749	install_requires=requirements,
750		750
751	extras_require={	751	extras_require={
752	},	752	},
753	package_data={	753	package_data={
754	'torch_npu': [	754	'torch_npu': [
755	'*.so',	755	'*.so',
756	'lib/.so',	756	'lib/.so',
757	],	757	],
758	'torchnpugen': [	758	'torchnpugen': [
759	'.py', '/.py',	759	'.py', '/.py',
760	'.yaml', '/.yaml',	760	'.yaml', '/.yaml',
761	'.json', '/.json',	761	'.json', '/.json',
762	'.cpp', '/.cpp',	762	'.cpp', '/.cpp',
763	'.h', '/.h',	763	'.h', '/.h',
764	],	764	],
765	},	765	},
766	cmdclass={	766	cmdclass={
767	'build_clib': CPPLibBuild,	767	'build_clib': CPPLibBuild,
768	'build_ext': Build,	768	'build_ext': Build,
769	'build_py': PythonPackageBuild,	769	'build_py': PythonPackageBuild,
770	'bdist_wheel': BdistWheelBuild,	770	'bdist_wheel': BdistWheelBuild,
771	'install': InstallCmd,	771	'install': InstallCmd,
772	'clean': Clean	772	'clean': Clean
773	},	773	},
774	entry_points={	774	entry_points={
775	'console_scripts': [	775	'console_scripts': [
776	'torch_npu_run = torch_npu.distributed.run:_main',	776	'torch_npu_run = torch_npu.distributed.run:_main',
777	],	777	],
778	'torch.backends': [	778	'torch.backends': [
779	'torch_npu = torch_npu:_autoload',	779	'torch_npu = torch_npu:_autoload',
780	],	780	],
781	}	781	}
782	)	782	)

Mtest/_inductor/test_empty.py+0-1

	@@ -42,4 +42,3 @@ instantiate_parametrized_tests(TestEmpty)
42		42
43	if __name__ == "__main__":	43	if __name__ == "__main__":
44	run_tests()	44	run_tests()
45

Mtest/_inductor/test_exp.py+0-1

	@@ -25,4 +25,3 @@ instantiate_parametrized_tests(TestExp)
25		25
26	if __name__ == "__main__":	26	if __name__ == "__main__":
27	run_tests()	27	run_tests()
28

Mtest/_inductor/test_fxpass_level.py+91-91

Mtest/_inductor/test_gt.py+0-1

	@@ -29,4 +29,3 @@ instantiate_parametrized_tests(TestGt)
29		29
30	if __name__ == "__main__":	30	if __name__ == "__main__":
31	run_tests()	31	run_tests()
32

Mtest/_inductor/test_shape_handling.py+815-815

	@@ -1,816 +1,816 @@
1	import unittest	1	import unittest
2	from unittest import mock	2	from unittest import mock
3		3
4	import torch	4	import torch
5	from torch.nn.parallel import scatter_gather	5	from torch.nn.parallel import scatter_gather
6	from torch.testing._internal.common_utils import (	6	from torch.testing._internal.common_utils import (
7	IS_JETSON,	7	IS_JETSON,
8	IS_REMOTE_GPU,	8	IS_REMOTE_GPU,
9	IS_SANDCASTLE,	9	IS_SANDCASTLE,
10	NoTest,	10	NoTest,
11	TEST_PRIVATEUSE1,	11	TEST_PRIVATEUSE1,
12	TestCase,	12	TestCase,
13	instantiate_parametrized_tests,	13	instantiate_parametrized_tests,
14	run_tests,	14	run_tests,
15	skipCUDANonDefaultStreamIf,	15	skipCUDANonDefaultStreamIf,
16	skipIfRocm,	16	skipIfRocm,
17	)	17	)
18	import torch_npu	18	import torch_npu
19	import torch_npu._inductor	19	import torch_npu._inductor
20	import torch_npu.testing	20	import torch_npu.testing
21	from torch_npu.testing.common_utils import get_cycles_per_ms	21	from torch_npu.testing.common_utils import get_cycles_per_ms
22	from torch_npu._inductor.shape_handling import unified_copy	22	from torch_npu._inductor.shape_handling import unified_copy
23	import torch_npu._inductor.shape_handling as shape_handling_module	23	import torch_npu._inductor.shape_handling as shape_handling_module
24		24
25	torch._dynamo.config.cache_size_limit = 128	25	torch._dynamo.config.cache_size_limit = 128
26		26
27	if not torch.npu.is_available():	27	if not torch.npu.is_available():
28	raise unittest.SkipTest("NPU is not available")	28	raise unittest.SkipTest("NPU is not available")
29		29
30	device = "npu"	30	device = "npu"
31		31
32		32
33	def model_fn(A, B):	33	def model_fn(A, B):
34	return A + B	34	return A + B
35		35
36	shape_options = {	36	shape_options = {
37	"enable_shape_handling": True,	37	"enable_shape_handling": True,
38	"shape_handling_configs": [	38	"shape_handling_configs": [
39	{	39	{
40	"type": "BATCHSIZE", # 处理的维度类型(Required)	40	"type": "BATCHSIZE", # 处理的维度类型(Required)
41	"dimensions": 0, # 该维度在tensor中的下标(BATCHSIZE默认为0)(Optional)	41	"dimensions": 0, # 该维度在tensor中的下标(BATCHSIZE默认为0)(Optional)
42	# "indices": [0, 1], # 需要处理的tensor下标(默认为所有tensor)(Optional)	42	# "indices": [0, 1], # 需要处理的tensor下标(默认为所有tensor)(Optional)
43	"value": 0.0, # padding时填充的值, 默认为0.0(Optional)	43	"value": 0.0, # padding时填充的值, 默认为0.0(Optional)
44	"gears": [], # 自定义档位信息(Optional)	44	"gears": [], # 自定义档位信息(Optional)
45	"min_size": 1, # 该维度的最小大小(档位), 默认为1(Optional)	45	"min_size": 1, # 该维度的最小大小(档位), 默认为1(Optional)
46	"max_size": 1024, # 该维度的最大大小(档位), 默认为1024(Optional)	46	"max_size": 1024, # 该维度的最大大小(档位), 默认为1024(Optional)
47	"policy": "TIMES", # 依据min_size, max_size自动生成gears的策略, 默认为TIMES, 表示生成范围内2的整数幂档位(Optional)	47	"policy": "TIMES", # 依据min_size, max_size自动生成gears的策略, 默认为TIMES, 表示生成范围内2的整数幂档位(Optional)
48	},	48	},
49	{	49	{
50	"type": "SEQLEN", # 处理的维度类型(Required)	50	"type": "SEQLEN", # 处理的维度类型(Required)
51	"dimensions": [1, 1], # 该维度在tensor中的下标, 与indices一一对应, 也可以接收dimension表示所有tensor使用相同的dimension index, 优先接收dimension(SEQLEN默认为1)(Optional)	51	"dimensions": [1, 1], # 该维度在tensor中的下标, 与indices一一对应, 也可以接收dimension表示所有tensor使用相同的dimension index, 优先接收dimension(SEQLEN默认为1)(Optional)
52	# "indices": [0, 1], # 需要处理的tensor下标(默认为所有tensor)(Optional)	52	# "indices": [0, 1], # 需要处理的tensor下标(默认为所有tensor)(Optional)
53	"value": 0.0, # padding时填充的值, 默认为0.0(Optional)	53	"value": 0.0, # padding时填充的值, 默认为0.0(Optional)
54	"gears": [], # 自定义档位信息(Optional)	54	"gears": [], # 自定义档位信息(Optional)
55	"min_size": 1, # 该维度的最小大小(档位), 默认为1(Optional)	55	"min_size": 1, # 该维度的最小大小(档位), 默认为1(Optional)
56	"max_size": 1024, # 该维度的最大大小(档位), 默认为1024(Optional)	56	"max_size": 1024, # 该维度的最大大小(档位), 默认为1024(Optional)
57	"policy": "TIMES", # 依据min_size, max_size自动生成gears的策略, 默认为TIMES, 表示生成范围内2的整数幂档位(Optional)	57	"policy": "TIMES", # 依据min_size, max_size自动生成gears的策略, 默认为TIMES, 表示生成范围内2的整数幂档位(Optional)
58	}	58	}
59	]	59	]
60	}	60	}
61		61
62		62
63	class TestShapeHandling(TestCase):	63	class TestShapeHandling(TestCase):
64	def test_init_no_input(self):	64	def test_init_no_input(self):
65	shape_handling = torch_npu._inductor.NPUShapeHandling()	65	shape_handling = torch_npu._inductor.NPUShapeHandling()
66	self.assertNotEqual(shape_handling, None)	66	self.assertNotEqual(shape_handling, None)
67		67
68	def test_init_with_empty_conifg(self):	68	def test_init_with_empty_conifg(self):
69	configs = []	69	configs = []
70	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	70	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
71	self.assertNotEqual(shape_handling, None)	71	self.assertNotEqual(shape_handling, None)
72		72
73	def test_init_with_gears(self):	73	def test_init_with_gears(self):
74	configs = [	74	configs = [
75	{	75	{
76	"type": "BATCHSIZE",	76	"type": "BATCHSIZE",
77	"gears": [16, 32, 64]	77	"gears": [16, 32, 64]
78	},	78	},
79	{	79	{
80	"type": "SEQLEN",	80	"type": "SEQLEN",
81	"gears": [16, 32, 64]	81	"gears": [16, 32, 64]
82	}	82	}
83	]	83	]
84	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	84	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
85	self.assertNotEqual(shape_handling, None)	85	self.assertNotEqual(shape_handling, None)
86		86
87	def test_init_with_policy(self):	87	def test_init_with_policy(self):
88	configs = [	88	configs = [
89	{	89	{
90	"type": "BATCHSIZE",	90	"type": "BATCHSIZE",
91	"min_size": 2,	91	"min_size": 2,
92	"max_size": 8,	92	"max_size": 8,
93	"policy": "TIMES"	93	"policy": "TIMES"
94	},	94	},
95	{	95	{
96	"type": "SEQLEN",	96	"type": "SEQLEN",
97	"min_size": 2,	97	"min_size": 2,
98	"max_size": 8,	98	"max_size": 8,
99	"policy": "TIMES"	99	"policy": "TIMES"
100	}	100	}
101	]	101	]
102	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	102	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
103	self.assertNotEqual(shape_handling, None)	103	self.assertNotEqual(shape_handling, None)
104		104
105	def test_transform_no_operation(self):	105	def test_transform_no_operation(self):
106	configs = [	106	configs = [
107	{	107	{
108	"type": "BATCHSIZE",	108	"type": "BATCHSIZE",
109	"gears": [32],	109	"gears": [32],
110	"dimensions": 0	110	"dimensions": 0
111	},	111	},
112	{	112	{
113	"type": "SEQLEN",	113	"type": "SEQLEN",
114	"gears": [128],	114	"gears": [128],
115	"dimensions": [1]	115	"dimensions": [1]
116	}	116	}
117	]	117	]
118	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	118	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
119	input_tensor = torch.randn(32, 128)	119	input_tensor = torch.randn(32, 128)
120	outputs = shape_handling.transform([input_tensor])	120	outputs = shape_handling.transform([input_tensor])
121	self.assertEqual(outputs[0][0].shape, input_tensor.shape)	121	self.assertEqual(outputs[0][0].shape, input_tensor.shape)
122		122
123	def test_transform_padding(self):	123	def test_transform_padding(self):
124	configs = [	124	configs = [
125	{	125	{
126	"type": "BATCHSIZE",	126	"type": "BATCHSIZE",
127	"min_size": 16,	127	"min_size": 16,
128	"max_size": 256,	128	"max_size": 256,
129	"policy": "TIMES",	129	"policy": "TIMES",
130	"dimensions": 0	130	"dimensions": 0
131	},	131	},
132	{	132	{
133	"type": "SEQLEN",	133	"type": "SEQLEN",
134	"min_size": 16,	134	"min_size": 16,
135	"max_size": 256,	135	"max_size": 256,
136	"policy": "TIMES",	136	"policy": "TIMES",
137	"dimensions": [1]	137	"dimensions": [1]
138	}	138	}
139	]	139	]
140	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	140	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
141	input_tensor = torch.randn(48, 96)	141	input_tensor = torch.randn(48, 96)
142	outputs = shape_handling.transform([input_tensor])	142	outputs = shape_handling.transform([input_tensor])
143	self.assertEqual(outputs[0][0].shape, (64, 128))	143	self.assertEqual(outputs[0][0].shape, (64, 128))
144		144
145	def test_transform_padding_with_indices(self):	145	def test_transform_padding_with_indices(self):
146	configs = [	146	configs = [
147	{	147	{
148	"type": "BATCHSIZE",	148	"type": "BATCHSIZE",
149	"min_size": 16,	149	"min_size": 16,
150	"max_size": 256,	150	"max_size": 256,
151	"policy": "TIMES",	151	"policy": "TIMES",
152	"dimensions": 0,	152	"dimensions": 0,
153	"indices": [0]	153	"indices": [0]
154	},	154	},
155	{	155	{
156	"type": "SEQLEN",	156	"type": "SEQLEN",
157	"min_size": 16,	157	"min_size": 16,
158	"max_size": 256,	158	"max_size": 256,
159	"policy": "TIMES",	159	"policy": "TIMES",
160	"dimensions": [1],	160	"dimensions": [1],
161	"indices": [0]	161	"indices": [0]
162	}	162	}
163	]	163	]
164	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	164	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
165	input_tensor1 = torch.randn(48, 96)	165	input_tensor1 = torch.randn(48, 96)
166	input_tensor2 = torch.randn(48, 96)	166	input_tensor2 = torch.randn(48, 96)
167	outputs = shape_handling.transform([input_tensor1, input_tensor2])	167	outputs = shape_handling.transform([input_tensor1, input_tensor2])
168	self.assertEqual(outputs[0][0].shape, (64, 128))	168	self.assertEqual(outputs[0][0].shape, (64, 128))
169	self.assertEqual(outputs[0][1].shape, (48, 96))	169	self.assertEqual(outputs[0][1].shape, (48, 96))
170		170
171	def test_transform_split(self):	171	def test_transform_split(self):
172	"""测试分割操作"""	172	"""测试分割操作"""
173	# 设置max_size为128，那么超过128的会被分割	173	# 设置max_size为128，那么超过128的会被分割
174	configs = [	174	configs = [
175	{	175	{
176	"type": "BATCHSIZE",	176	"type": "BATCHSIZE",
177	"gears": [64, 128],	177	"gears": [64, 128],
178	"dimensions": 0	178	"dimensions": 0
179	},	179	},
180	{	180	{
181	"type": "SEQLEN",	181	"type": "SEQLEN",
182	"gears": [64, 128],	182	"gears": [64, 128],
183	"dimensions": [1]	183	"dimensions": [1]
184	}	184	}
185	]	185	]
186	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	186	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
187	input_tensor = torch.randn(200, 96) # 超过max_size	187	input_tensor = torch.randn(200, 96) # 超过max_size
188	outputs = shape_handling.transform([input_tensor])	188	outputs = shape_handling.transform([input_tensor])
189	# 验证分割结果：分割为两个组，第一段128，第二段72，再填充为128	189	# 验证分割结果：分割为两个组，第一段128，第二段72，再填充为128
190	self.assertEqual(len(outputs), 2) # 分割为两个组	190	self.assertEqual(len(outputs), 2) # 分割为两个组
191	self.assertEqual(outputs[0][0].shape, (128, 128)) # 第一段	191	self.assertEqual(outputs[0][0].shape, (128, 128)) # 第一段
192	self.assertEqual(outputs[1][0].shape, (128, 128)) # 第二段	192	self.assertEqual(outputs[1][0].shape, (128, 128)) # 第二段
193		193
194	def test_recover_padding(self):	194	def test_recover_padding(self):
195	"""测试恢复填充的张量"""	195	"""测试恢复填充的张量"""
196	configs = [	196	configs = [
197	{	197	{
198	"type": "BATCHSIZE",	198	"type": "BATCHSIZE",
199	"min_size": 16,	199	"min_size": 16,
200	"max_size": 256,	200	"max_size": 256,
201	"policy": "TIMES",	201	"policy": "TIMES",
202	"dimensions": 0,	202	"dimensions": 0,
203	"indices": [0]	203	"indices": [0]
204	},	204	},
205	{	205	{
206	"type": "SEQLEN",	206	"type": "SEQLEN",
207	"min_size": 16,	207	"min_size": 16,
208	"max_size": 256,	208	"max_size": 256,
209	"policy": "TIMES",	209	"policy": "TIMES",
210	"dimensions": [1],	210	"dimensions": [1],
211	"indices": [0]	211	"indices": [0]
212	}	212	}
213	]	213	]
214	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	214	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
215	orig_tensor = torch.randn(48, 128)	215	orig_tensor = torch.randn(48, 128)
216	padded_group = shape_handling.transform([orig_tensor])	216	padded_group = shape_handling.transform([orig_tensor])
217	# 执行恢复	217	# 执行恢复
218	recovered = shape_handling.recover(padded_group)	218	recovered = shape_handling.recover(padded_group)
219	# 验证尺寸还原	219	# 验证尺寸还原
220	self.assertEqual(recovered[0].shape, orig_tensor.shape)	220	self.assertEqual(recovered[0].shape, orig_tensor.shape)
221	# 验证数据一致性	221	# 验证数据一致性
222	self.assertTrue(torch.allclose(recovered[0], orig_tensor))	222	self.assertTrue(torch.allclose(recovered[0], orig_tensor))
223		223
224	def test_recover_split(self):	224	def test_recover_split(self):
225	"""测试恢复分割的张量组"""	225	"""测试恢复分割的张量组"""
226	configs = [	226	configs = [
227	{	227	{
228	"type": "BATCHSIZE",	228	"type": "BATCHSIZE",
229	"min_size": 16,	229	"min_size": 16,
230	"max_size": 128,	230	"max_size": 128,
231	"policy": "TIMES",	231	"policy": "TIMES",
232	"dimensions": 0,	232	"dimensions": 0,
233	"indices": [0]	233	"indices": [0]
234	},	234	},
235	{	235	{
236	"type": "SEQLEN",	236	"type": "SEQLEN",
237	"min_size": 16,	237	"min_size": 16,
238	"max_size": 256,	238	"max_size": 256,
239	"policy": "TIMES",	239	"policy": "TIMES",
240	"dimensions": [1],	240	"dimensions": [1],
241	"indices": [0]	241	"indices": [0]
242	}	242	}
243	]	243	]
244	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	244	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
245	orig_tensor = torch.randn(200, 128)	245	orig_tensor = torch.randn(200, 128)
246	split_groups = shape_handling.transform([orig_tensor])	246	split_groups = shape_handling.transform([orig_tensor])
247	# 执行恢复	247	# 执行恢复
248	recovered = shape_handling.recover(split_groups)	248	recovered = shape_handling.recover(split_groups)
249	# 验证拼接还原	249	# 验证拼接还原
250	self.assertEqual(recovered[0].shape, orig_tensor.shape)	250	self.assertEqual(recovered[0].shape, orig_tensor.shape)
251	self.assertTrue(torch.allclose(recovered[0], orig_tensor))	251	self.assertTrue(torch.allclose(recovered[0], orig_tensor))
252		252
253	def test_invalid_dim(self):	253	def test_invalid_dim(self):
254	"""测试无效维度异常"""	254	"""测试无效维度异常"""
255	configs = [	255	configs = [
256	{	256	{
257	"type": "BATCHSIZE",	257	"type": "BATCHSIZE",
258	"min_size": 16,	258	"min_size": 16,
259	"max_size": 128,	259	"max_size": 128,
260	"policy": "TIMES",	260	"policy": "TIMES",
261	"dimensions": 3,	261	"dimensions": 3,
262	"indices": [0]	262	"indices": [0]
263	}	263	}
264	]	264	]
265	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	265	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
266	input_tensor = torch.randn(32, 128)	266	input_tensor = torch.randn(32, 128)
267	# 验证越界维度触发异常	267	# 验证越界维度触发异常
268	with self.assertRaises(RuntimeError):	268	with self.assertRaises(RuntimeError):
269	shape_handling.transform([input_tensor]) # 有效维度应为0或1	269	shape_handling.transform([input_tensor]) # 有效维度应为0或1
270		270
271	def test_register_custom_strategy(self):	271	def test_register_custom_strategy(self):
272	"""测试注册自定义策略"""	272	"""测试注册自定义策略"""
273	# 创建自定义策略类	273	# 创建自定义策略类
274	class CustomBsShapeOp(torch_npu._C._BSShapeOpStrategy):	274	class CustomBsShapeOp(torch_npu._C._BSShapeOpStrategy):
275	def __init__(self):	275	def __init__(self):
276	super().__init__()	276	super().__init__()
277	self.transform_called = False	277	self.transform_called = False
278	self.recover_called = False	278	self.recover_called = False
279		279
280	def Transform(self, inputs, outputs):	280	def Transform(self, inputs, outputs):
281	self.transform_called = True	281	self.transform_called = True
282	# 简单实现：直接复制输入到输出	282	# 简单实现：直接复制输入到输出
283	if inputs:	283	if inputs:
284	outputs.append([tensor.clone() for tensor in inputs])	284	outputs.append([tensor.clone() for tensor in inputs])
285		285
286	def Recover(self, inputs, outputs):	286	def Recover(self, inputs, outputs):
287	self.recover_called = True	287	self.recover_called = True
288	# 简单实现：直接复制第一个组的第一个张量	288	# 简单实现：直接复制第一个组的第一个张量
289	if inputs and inputs[0]:	289	if inputs and inputs[0]:
290	outputs.append(inputs[0][0].clone())	290	outputs.append(inputs[0][0].clone())
291		291
292	class CustomSeqShapeOp(torch_npu._C._SeqShapeOpStrategy):	292	class CustomSeqShapeOp(torch_npu._C._SeqShapeOpStrategy):
293	def __init__(self):	293	def __init__(self):
294	super().__init__()	294	super().__init__()
295	self.transform_called = False	295	self.transform_called = False
296	self.recover_called = False	296	self.recover_called = False
297		297
298	def Transform(self, inputs, outputs):	298	def Transform(self, inputs, outputs):
299	self.transform_called = True	299	self.transform_called = True
300	# 简单实现：直接复制输入到输出	300	# 简单实现：直接复制输入到输出
301	if inputs:	301	if inputs:
302	outputs.append([tensor.clone() for tensor in inputs])	302	outputs.append([tensor.clone() for tensor in inputs])
303		303
304	configs = [	304	configs = [
305	{	305	{
306	"type": "BATCHSIZE",	306	"type": "BATCHSIZE",
307	"min_size": 16,	307	"min_size": 16,
308	"max_size": 128,	308	"max_size": 128,
309	"policy": "TIMES",	309	"policy": "TIMES",
310	"dimensions": 0,	310	"dimensions": 0,
311	"indices": [0]	311	"indices": [0]
312	},	312	},
313	{	313	{
314	"type": "SEQLEN",	314	"type": "SEQLEN",
315	"min_size": 16,	315	"min_size": 16,
316	"max_size": 256,	316	"max_size": 256,
317	"policy": "TIMES",	317	"policy": "TIMES",
318	"dimensions": [1],	318	"dimensions": [1],
319	"indices": [0]	319	"indices": [0]
320	}	320	}
321	]	321	]
322	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	322	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
323	custom_bs_strategy = CustomBsShapeOp()	323	custom_bs_strategy = CustomBsShapeOp()
324	custom_seq_strategy = CustomSeqShapeOp()	324	custom_seq_strategy = CustomSeqShapeOp()
325		325
326	# 注册自定义策略	326	# 注册自定义策略
327	shape_handling.register_batch_size_strategy(custom_bs_strategy)	327	shape_handling.register_batch_size_strategy(custom_bs_strategy)
328	shape_handling.register_sequence_strategy(custom_seq_strategy)	328	shape_handling.register_sequence_strategy(custom_seq_strategy)
329		329
330	# 验证自定义策略被调用	330	# 验证自定义策略被调用
331	input_tensor = torch.randn(32, 128)	331	input_tensor = torch.randn(32, 128)
332	shape_handling.transform([input_tensor])	332	shape_handling.transform([input_tensor])
333		333
334	self.assertTrue(custom_bs_strategy.transform_called)	334	self.assertTrue(custom_bs_strategy.transform_called)
335	self.assertTrue(custom_seq_strategy.transform_called)	335	self.assertTrue(custom_seq_strategy.transform_called)
336		336
337		337
338		338
339	class TestShapeHandlingBranchCoverage(TestCase):	339	class TestShapeHandlingBranchCoverage(TestCase):
340	def test_validate_configs_error_branches(self):	340	def test_validate_configs_error_branches(self):
341	shape_handling = torch_npu._inductor.NPUShapeHandling()	341	shape_handling = torch_npu._inductor.NPUShapeHandling()
342		342
343	with self.assertRaises(ValueError):	343	with self.assertRaises(ValueError):
344	shape_handling._validate_configs([{}])	344	shape_handling._validate_configs([{}])
345		345
346	with self.assertRaises(ValueError):	346	with self.assertRaises(ValueError):
347	shape_handling._validate_configs([{"type": 1}])	347	shape_handling._validate_configs([{"type": 1}])
348		348
349	with self.assertRaises(ValueError):	349	with self.assertRaises(ValueError):
350	shape_handling._validate_configs([{"type": "UNKNOWN"}])	350	shape_handling._validate_configs([{"type": "UNKNOWN"}])
351		351
352	with self.assertRaises(ValueError):	352	with self.assertRaises(ValueError):
353	shape_handling._validate_configs([{"type": "BATCHSIZE", "dimensions": "0"}])	353	shape_handling._validate_configs([{"type": "BATCHSIZE", "dimensions": "0"}])
354		354
355	with self.assertRaises(ValueError):	355	with self.assertRaises(ValueError):
356	shape_handling._validate_configs([{"type": "BATCHSIZE", "dimensions": [0, "1"]}])	356	shape_handling._validate_configs([{"type": "BATCHSIZE", "dimensions": [0, "1"]}])
357		357
358	with self.assertRaises(ValueError):	358	with self.assertRaises(ValueError):
359	shape_handling._validate_configs([{"type": "BATCHSIZE", "min_size": 1.5}])	359	shape_handling._validate_configs([{"type": "BATCHSIZE", "min_size": 1.5}])
360		360
361	with self.assertRaises(ValueError):	361	with self.assertRaises(ValueError):
362	shape_handling._validate_configs([{"type": "BATCHSIZE", "value": "bad"}])	362	shape_handling._validate_configs([{"type": "BATCHSIZE", "value": "bad"}])
363		363
364	with self.assertRaises(ValueError):	364	with self.assertRaises(ValueError):
365	shape_handling._validate_configs([{"type": "BATCHSIZE", "policy": 1}])	365	shape_handling._validate_configs([{"type": "BATCHSIZE", "policy": 1}])
366		366
367	with self.assertRaises(ValueError):	367	with self.assertRaises(ValueError):
368	shape_handling._validate_configs([{"type": "BATCHSIZE", "policy": "BAD"}])	368	shape_handling._validate_configs([{"type": "BATCHSIZE", "policy": "BAD"}])
369		369
370	with self.assertRaises(ValueError):	370	with self.assertRaises(ValueError):
371	shape_handling._validate_configs([{"type": "BATCHSIZE"}, {"type": "BATCHSIZE"}])	371	shape_handling._validate_configs([{"type": "BATCHSIZE"}, {"type": "BATCHSIZE"}])
372		372
373	with self.assertRaises(ValueError):	373	with self.assertRaises(ValueError):
374	shape_handling._validate_configs(	374	shape_handling._validate_configs(
375	[{"type": "BATCHSIZE"}, {"type": "SEQLEN"}, {"type": "BATCHSIZE"}]	375	[{"type": "BATCHSIZE"}, {"type": "SEQLEN"}, {"type": "BATCHSIZE"}]
376	)	376	)
377		377
378	def test_construct_indices_branches(self):	378	def test_construct_indices_branches(self):
379	shape_handling = torch_npu._inductor.NPUShapeHandling()	379	shape_handling = torch_npu._inductor.NPUShapeHandling()
380	tensors = [torch.randn(2, 3), torch.randn(2)]	380	tensors = [torch.randn(2, 3), torch.randn(2)]
381		381
382	bs_indices = shape_handling._construct_indices(tensors, [], "BATCHSIZE")	382	bs_indices = shape_handling._construct_indices(tensors, [], "BATCHSIZE")
383	self.assertEqual(bs_indices, [0, 1])	383	self.assertEqual(bs_indices, [0, 1])
384		384
385	seq_indices_default = shape_handling._construct_indices(tensors, [], "SEQLEN")	385	seq_indices_default = shape_handling._construct_indices(tensors, [], "SEQLEN")
386	self.assertEqual(seq_indices_default, [0])	386	self.assertEqual(seq_indices_default, [0])
387		387
388	seq_indices_dim0 = shape_handling._construct_indices(tensors, [0], "SEQLEN")	388	seq_indices_dim0 = shape_handling._construct_indices(tensors, [0], "SEQLEN")
389	self.assertEqual(seq_indices_dim0, [0, 1])	389	self.assertEqual(seq_indices_dim0, [0, 1])
390		390
391	def test_delay_initialize_builds_missing_indices(self):	391	def test_delay_initialize_builds_missing_indices(self):
392	configs = [	392	configs = [
393	{	393	{
394	"type": "BATCHSIZE",	394	"type": "BATCHSIZE",
395	"dimensions": [0],	395	"dimensions": [0],
396	"indices": [],	396	"indices": [],
397	"gears": [8]	397	"gears": [8]
398	},	398	},
399	{	399	{
400	"type": "SEQLEN",	400	"type": "SEQLEN",
401	"dimensions": [1],	401	"dimensions": [1],
402	"indices": [0],	402	"indices": [0],
403	"gears": [8]	403	"gears": [8]
404	}	404	}
405	]	405	]
406	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)	406	shape_handling = torch_npu._inductor.NPUShapeHandling(configs)
407	self.assertTrue(shape_handling.delay_init)	407	self.assertTrue(shape_handling.delay_init)
408		408
409	shape_handling.delay_initialize([torch.randn(4, 5)])	409	shape_handling.delay_initialize([torch.randn(4, 5)])
410	self.assertEqual(shape_handling.configs[0]["indices"], [0])	410	self.assertEqual(shape_handling.configs[0]["indices"], [0])
411	self.assertEqual(shape_handling.configs[1]["indices"], [0])	411	self.assertEqual(shape_handling.configs[1]["indices"], [0])
412	self.assertFalse(shape_handling.delay_init)	412	self.assertFalse(shape_handling.delay_init)
413		413
414	def test_flatten_and_unflatten_with_and_without_tensor(self):	414	def test_flatten_and_unflatten_with_and_without_tensor(self):
415	shape_handling = torch_npu._inductor.NPUShapeHandling()	415	shape_handling = torch_npu._inductor.NPUShapeHandling()
416		416
417	tensors, indices, leaves, spec = shape_handling.flatten_to_tensors({"a": 1, "b": "x"})	417	tensors, indices, leaves, spec = shape_handling.flatten_to_tensors({"a": 1, "b": "x"})
418	self.assertEqual(list(tensors), [])	418	self.assertEqual(list(tensors), [])
419	self.assertEqual(list(indices), [])	419	self.assertEqual(list(indices), [])
420	self.assertEqual(len(leaves), 2)	420	self.assertEqual(len(leaves), 2)
421	self.assertIsNotNone(spec)	421	self.assertIsNotNone(spec)
422		422
423	structure = ((torch.tensor([1.0]), "k"), {"v": torch.tensor([2.0])})	423	structure = ((torch.tensor([1.0]), "k"), {"v": torch.tensor([2.0])})
424	tensors, indices, leaves, spec = shape_handling.flatten_to_tensors(structure)	424	tensors, indices, leaves, spec = shape_handling.flatten_to_tensors(structure)
425	rebuilt = shape_handling.unflatten_from_tensors(	425	rebuilt = shape_handling.unflatten_from_tensors(
426	[torch.tensor([10.0]), torch.tensor([20.0])], indices, leaves, spec	426	[torch.tensor([10.0]), torch.tensor([20.0])], indices, leaves, spec
427	)	427	)
428	self.assertEqual(rebuilt[0][1], "k")	428	self.assertEqual(rebuilt[0][1], "k")
429	self.assertTrue(torch.equal(rebuilt[0][0], torch.tensor([10.0])))	429	self.assertTrue(torch.equal(rebuilt[0][0], torch.tensor([10.0])))
430	self.assertTrue(torch.equal(rebuilt[1]["v"], torch.tensor([20.0])))	430	self.assertTrue(torch.equal(rebuilt[1]["v"], torch.tensor([20.0])))
431		431
432	def test_transform_hook_default_and_custom_paths(self):	432	def test_transform_hook_default_and_custom_paths(self):
433	shape_handling = torch_npu._inductor.NPUShapeHandling()	433	shape_handling = torch_npu._inductor.NPUShapeHandling()
434	input_a = torch.tensor([1.0])	434	input_a = torch.tensor([1.0])
435	input_b = torch.tensor([2.0])	435	input_b = torch.tensor([2.0])
436	transform_res = [	436	transform_res = [
437	[torch.tensor([3.0]), torch.tensor([4.0])],	437	[torch.tensor([3.0]), torch.tensor([4.0])],
438	[torch.tensor([5.0]), torch.tensor([6.0])]	438	[torch.tensor([5.0]), torch.tensor([6.0])]
439	]	439	]
440		440
441	with mock.patch.object(	441	with mock.patch.object(
442	torch_npu._inductor.NPUShapeHandling, "transform", return_value=transform_res	442	torch_npu._inductor.NPUShapeHandling, "transform", return_value=transform_res
443	):	443	):
444	args_list, kwargs_list = shape_handling.transform_hook(input_a, extra=input_b)	444	args_list, kwargs_list = shape_handling.transform_hook(input_a, extra=input_b)
445	self.assertEqual(len(args_list), 2)	445	self.assertEqual(len(args_list), 2)
446	self.assertEqual(len(kwargs_list), 2)	446	self.assertEqual(len(kwargs_list), 2)
447	self.assertTrue(torch.equal(args_list[0][0], torch.tensor([3.0])))	447	self.assertTrue(torch.equal(args_list[0][0], torch.tensor([3.0])))
448	self.assertTrue(torch.equal(kwargs_list[1]["extra"], torch.tensor([6.0])))	448	self.assertTrue(torch.equal(kwargs_list[1]["extra"], torch.tensor([6.0])))
449		449
450	recorded = {}	450	recorded = {}
451		451
452	def pre_fn(args, *kwargs):	452	def pre_fn(args, *kwargs):
453	recorded["pre"] = (args, kwargs)	453	recorded["pre"] = (args, kwargs)
454	return [args[0]]	454	return [args[0]]
455		455
456	def post_fn(outputs):	456	def post_fn(outputs):
457	recorded["post"] = outputs	457	recorded["post"] = outputs
458	return [("ok",)], [{"done": True}]	458	return [("ok",)], [{"done": True}]
459		459
460	shape_handling_custom = torch_npu._inductor.NPUShapeHandling(	460	shape_handling_custom = torch_npu._inductor.NPUShapeHandling(
461	transform_pre_fn=pre_fn,	461	transform_pre_fn=pre_fn,
462	transform_post_fn=post_fn,	462	transform_post_fn=post_fn,
463	)	463	)
464	with mock.patch.object(	464	with mock.patch.object(
465	torch_npu._inductor.NPUShapeHandling, "transform", return_value=[[torch.tensor([9.0])]]	465	torch_npu._inductor.NPUShapeHandling, "transform", return_value=[[torch.tensor([9.0])]]
466	):	466	):
467	out_args, out_kwargs = shape_handling_custom.transform_hook(torch.tensor([7.0]))	467	out_args, out_kwargs = shape_handling_custom.transform_hook(torch.tensor([7.0]))
468	self.assertIn("pre", recorded)	468	self.assertIn("pre", recorded)
469	self.assertIn("post", recorded)	469	self.assertIn("post", recorded)
470	self.assertEqual(out_args, [("ok",)])	470	self.assertEqual(out_args, [("ok",)])
471	self.assertEqual(out_kwargs, [{"done": True}])	471	self.assertEqual(out_kwargs, [{"done": True}])
472		472
473	shape_handling_none = torch_npu._inductor.NPUShapeHandling(transform_post_fn=lambda _: None)	473	shape_handling_none = torch_npu._inductor.NPUShapeHandling(transform_post_fn=lambda _: None)
474	with mock.patch.object(	474	with mock.patch.object(
475	torch_npu._inductor.NPUShapeHandling, "transform", return_value=[[torch.tensor([1.0])]]	475	torch_npu._inductor.NPUShapeHandling, "transform", return_value=[[torch.tensor([1.0])]]
476	):	476	):
477	self.assertIsNone(shape_handling_none.transform_hook(torch.tensor([1.0])))	477	self.assertIsNone(shape_handling_none.transform_hook(torch.tensor([1.0])))
478		478
479	def test_recover_hook_default_and_custom_paths(self):	479	def test_recover_hook_default_and_custom_paths(self):
480	shape_handling = torch_npu._inductor.NPUShapeHandling()	480	shape_handling = torch_npu._inductor.NPUShapeHandling()
481	groups = [	481	groups = [
482	((torch.tensor([1.0]), "x"), {"y": torch.tensor([2.0])})	482	((torch.tensor([1.0]), "x"), {"y": torch.tensor([2.0])})
483	]	483	]
484		484
485	with mock.patch.object(	485	with mock.patch.object(
486	torch_npu._inductor.NPUShapeHandling,	486	torch_npu._inductor.NPUShapeHandling,
487	"recover",	487	"recover",
488	side_effect=lambda tensor_groups: list(tensor_groups[0]),	488	side_effect=lambda tensor_groups: list(tensor_groups[0]),
489	):	489	):
490	outputs = shape_handling.recover_hook(groups)	490	outputs = shape_handling.recover_hook(groups)
491	self.assertEqual(outputs[0][1], "x")	491	self.assertEqual(outputs[0][1], "x")
492	self.assertTrue(torch.equal(outputs[0][0], torch.tensor([1.0])))	492	self.assertTrue(torch.equal(outputs[0][0], torch.tensor([1.0])))
493	self.assertTrue(torch.equal(outputs[1]["y"], torch.tensor([2.0])))	493	self.assertTrue(torch.equal(outputs[1]["y"], torch.tensor([2.0])))
494		494
495	custom = torch_npu._inductor.NPUShapeHandling(	495	custom = torch_npu._inductor.NPUShapeHandling(
496	recover_pre_fn=lambda groups: [[torch.tensor([5.0])]],	496	recover_pre_fn=lambda groups: [[torch.tensor([5.0])]],
497	recover_post_fn=lambda recover_res: {"result": recover_res},	497	recover_post_fn=lambda recover_res: {"result": recover_res},
498	)	498	)
499	with mock.patch.object(	499	with mock.patch.object(
500	torch_npu._inductor.NPUShapeHandling, "recover", return_value=[torch.tensor([8.0])]	500	torch_npu._inductor.NPUShapeHandling, "recover", return_value=[torch.tensor([8.0])]
501	):	501	):
502	outputs = custom.recover_hook(groups)	502	outputs = custom.recover_hook(groups)
503	self.assertEqual(list(outputs.keys()), ["result"])	503	self.assertEqual(list(outputs.keys()), ["result"])
504	self.assertTrue(torch.equal(outputs["result"][0], torch.tensor([8.0])))	504	self.assertTrue(torch.equal(outputs["result"][0], torch.tensor([8.0])))
505		505
506	def test_get_shape_safe_and_patch_shape_handling(self):	506	def test_get_shape_safe_and_patch_shape_handling(self):
507	shape_handling = torch_npu._inductor.NPUShapeHandling()	507	shape_handling = torch_npu._inductor.NPUShapeHandling()
508	shape_info = shape_handling.get_shape_safe((torch.randn(2, 3), [torch.randn(1)]))	508	shape_info = shape_handling.get_shape_safe((torch.randn(2, 3), [torch.randn(1)]))
509	self.assertEqual(shape_info[0], [2, 3])	509	self.assertEqual(shape_info[0], [2, 3])
510	self.assertEqual(shape_info[1][0], [1])	510	self.assertEqual(shape_info[1][0], [1])
511	self.assertIs(shape_handling.get_shape_safe(1), int)	511	self.assertIs(shape_handling.get_shape_safe(1), int)
512		512
513	if hasattr(shape_handling_module.patch_shape_handling, "_is_patched"):	513	if hasattr(shape_handling_module.patch_shape_handling, "_is_patched"):
514	delattr(shape_handling_module.patch_shape_handling, "_is_patched")	514	delattr(shape_handling_module.patch_shape_handling, "_is_patched")
515	called = []	515	called = []
516	with mock.patch.object(shape_handling_module, "patch_dynamo_context", side_effect=lambda: called.append(1)):	516	with mock.patch.object(shape_handling_module, "patch_dynamo_context", side_effect=lambda: called.append(1)):
517	shape_handling_module.patch_shape_handling()	517	shape_handling_module.patch_shape_handling()
518	shape_handling_module.patch_shape_handling()	518	shape_handling_module.patch_shape_handling()
519	self.assertEqual(len(called), 1)	519	self.assertEqual(len(called), 1)
520	if hasattr(shape_handling_module.patch_shape_handling, "_is_patched"):	520	if hasattr(shape_handling_module.patch_shape_handling, "_is_patched"):
521	delattr(shape_handling_module.patch_shape_handling, "_is_patched")	521	delattr(shape_handling_module.patch_shape_handling, "_is_patched")
522		522
523		523
524	class TestDynamicShapeCompile(TestCase):	524	class TestDynamicShapeCompile(TestCase):
525	def test_npu_dynamic_shape_reuse_with_no_bucket(self):	525	def test_npu_dynamic_shape_reuse_with_no_bucket(self):
526		526
527	compiled_fn = torch.compile(	527	compiled_fn = torch.compile(
528	model_fn,	528	model_fn,
529	backend='inductor',	529	backend='inductor',
530	dynamic=False,	530	dynamic=False,
531	)	531	)
532		532
533	# 运行不同形状，验证是否只触发 4 次编译	533	# 运行不同形状，验证是否只触发 4 次编译
534	test_shapes = [(3, 20), (4, 20), (5, 20), (6, 20)]	534	test_shapes = [(3, 20), (4, 20), (5, 20), (6, 20)]
535		535
536		536
537	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):	537	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):
538	torch._inductor.metrics.generated_kernel_count.reset()	538	torch._inductor.metrics.generated_kernel_count.reset()
539	else:	539	else:
540	torch._inductor.metrics.generated_kernel_count = 0	540	torch._inductor.metrics.generated_kernel_count = 0
541		541
542	for shape in test_shapes:	542	for shape in test_shapes:
543	A = torch.randn(shape, device=device)	543	A = torch.randn(shape, device=device)
544	B = torch.randn(shape, device=device)	544	B = torch.randn(shape, device=device)
545	out = compiled_fn(A, B)	545	out = compiled_fn(A, B)
546	self.assertTrue(torch.allclose(out, A + B))	546	self.assertTrue(torch.allclose(out, A + B))
547		547
548		548
549	compile_count = torch._inductor.metrics.generated_kernel_count	549	compile_count = torch._inductor.metrics.generated_kernel_count
550		550
551	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）	551	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）
552	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")	552	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")
553		553
554	# 如果动态形状分档不生效，编译次数应为 4	554	# 如果动态形状分档不生效，编译次数应为 4
555	self.assertEqual(compile_count, 4)	555	self.assertEqual(compile_count, 4)
556		556
557	def test_npu_dynamic_shape_reuse_with_bucket(self):	557	def test_npu_dynamic_shape_reuse_with_bucket(self):
558		558
559	compiled_fn = torch.compile(	559	compiled_fn = torch.compile(
560	model_fn,	560	model_fn,
561	backend='inductor',	561	backend='inductor',
562	dynamic=False,	562	dynamic=False,
563	options=shape_options	563	options=shape_options
564	)	564	)
565		565
566	# 运行不同形状，验证是否只触发 2 次编译	566	# 运行不同形状，验证是否只触发 2 次编译
567	test_shapes = [(3, 32), (4, 32), (5, 32), (6, 32)]	567	test_shapes = [(3, 32), (4, 32), (5, 32), (6, 32)]
568	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):	568	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):
569	torch._inductor.metrics.generated_kernel_count.reset()	569	torch._inductor.metrics.generated_kernel_count.reset()
570	else:	570	else:
571	torch._inductor.metrics.generated_kernel_count = 0	571	torch._inductor.metrics.generated_kernel_count = 0
572		572
573	for shape in test_shapes:	573	for shape in test_shapes:
574	A = torch.randn(shape, device=device)	574	A = torch.randn(shape, device=device)
575	B = torch.randn(shape, device=device)	575	B = torch.randn(shape, device=device)
576	out = compiled_fn(A, B)	576	out = compiled_fn(A, B)
577	self.assertTrue(torch.allclose(out, A + B))	577	self.assertTrue(torch.allclose(out, A + B))
578		578
579		579
580	compile_count = torch._inductor.metrics.generated_kernel_count	580	compile_count = torch._inductor.metrics.generated_kernel_count
581		581
582	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）	582	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）
583	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")	583	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")
584		584
585	# 如果动态形状分档生效，编译次数应为 2	585	# 如果动态形状分档生效，编译次数应为 2
586	self.assertEqual(compile_count, 2)	586	self.assertEqual(compile_count, 2)
587		587
588	def test_npu_dynamic_shape_reuse_with_symbolic_shape(self):	588	def test_npu_dynamic_shape_reuse_with_symbolic_shape(self):
589		589
590	compiled_fn = torch.compile(	590	compiled_fn = torch.compile(
591	model_fn,	591	model_fn,
592	backend='inductor',	592	backend='inductor',
593	dynamic=True,	593	dynamic=True,
594	)	594	)
595		595
596	# 运行不同形状，验证是否只触发 1 次编译	596	# 运行不同形状，验证是否只触发 1 次编译
597	test_shapes = [(3, 32), (4, 32), (5, 32), (6, 32)]	597	test_shapes = [(3, 32), (4, 32), (5, 32), (6, 32)]
598		598
599	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):	599	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):
600	torch._inductor.metrics.generated_kernel_count.reset()	600	torch._inductor.metrics.generated_kernel_count.reset()
601	else:	601	else:
602	torch._inductor.metrics.generated_kernel_count = 0	602	torch._inductor.metrics.generated_kernel_count = 0
603		603
604	for shape in test_shapes:	604	for shape in test_shapes:
605	A = torch.randn(shape, device=device)	605	A = torch.randn(shape, device=device)
606	B = torch.randn(shape, device=device)	606	B = torch.randn(shape, device=device)
607	out = compiled_fn(A, B)	607	out = compiled_fn(A, B)
608	self.assertTrue(torch.allclose(out, A + B))	608	self.assertTrue(torch.allclose(out, A + B))
609		609
610	compile_count = torch._inductor.metrics.generated_kernel_count	610	compile_count = torch._inductor.metrics.generated_kernel_count
611		611
612		612
613		613
614	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）	614	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）
615	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")	615	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")
616		616
617	# 如果动态形状符号化生效，编译次数应为 1	617	# 如果动态形状符号化生效，编译次数应为 1
618	self.assertEqual(compile_count, 1)	618	self.assertEqual(compile_count, 1)
619		619
620	def test_npu_dynamic_shape_reuse_with_symbolic_shape_and_bucket(self):	620	def test_npu_dynamic_shape_reuse_with_symbolic_shape_and_bucket(self):
621		621
622	compiled_fn = torch.compile(	622	compiled_fn = torch.compile(
623	model_fn,	623	model_fn,
624	backend='inductor',	624	backend='inductor',
625	dynamic=True,	625	dynamic=True,
626	options=shape_options	626	options=shape_options
627	)	627	)
628		628
629	# 运行不同形状，验证是否只触发 1 次编译	629	# 运行不同形状，验证是否只触发 1 次编译
630	test_shapes = [(3, 32), (4, 32), (5, 32), (6, 32)]	630	test_shapes = [(3, 32), (4, 32), (5, 32), (6, 32)]
631		631
632	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):	632	if hasattr(torch._inductor.metrics.generated_kernel_count, 'reset'):
633	torch._inductor.metrics.generated_kernel_count.reset()	633	torch._inductor.metrics.generated_kernel_count.reset()
634	else:	634	else:
635	torch._inductor.metrics.generated_kernel_count = 0	635	torch._inductor.metrics.generated_kernel_count = 0
636		636
637	for shape in test_shapes:	637	for shape in test_shapes:
638	A = torch.randn(shape, device=device)	638	A = torch.randn(shape, device=device)
639	B = torch.randn(shape, device=device)	639	B = torch.randn(shape, device=device)
640	out = compiled_fn(A, B)	640	out = compiled_fn(A, B)
641	self.assertTrue(torch.allclose(out, A + B))	641	self.assertTrue(torch.allclose(out, A + B))
642		642
643	compile_count = torch._inductor.metrics.generated_kernel_count	643	compile_count = torch._inductor.metrics.generated_kernel_count
644		644
645		645
646	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）	646	# 获取 Inductor 编译的总次数（生成的 kernel 数量相关）
647	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")	647	print(f"\n[结果] Inductor 编译次数 (generated_kernel_count): {compile_count}")
648		648
649	# 如果动态形状符号化生效，编译次数应为 1	649	# 如果动态形状符号化生效，编译次数应为 1
650	self.assertEqual(compile_count, 1)	650	self.assertEqual(compile_count, 1)
651		651
652	def test_npu_invalid_shape_options_key_error(self):	652	def test_npu_invalid_shape_options_key_error(self):
653		653
654	# 构造一个非法的配置：最小尺寸大于最大尺寸	654	# 构造一个非法的配置：最小尺寸大于最大尺寸
655	invalid_shape_options = {	655	invalid_shape_options = {
656	"enable_shape_handling": True,	656	"enable_shape_handling": True,
657	"shape_handling_configs": [	657	"shape_handling_configs": [
658	{	658	{
659	"type": "BATCHSIZE",	659	"type": "BATCHSIZE",
660	"min_size": 1024, # 错误：min > max	660	"min_size": 1024, # 错误：min > max
661	"max_size": 1,	661	"max_size": 1,
662	"policy": "TIMES",	662	"policy": "TIMES",
663	}	663	}
664	]	664	]
665	}	665	}
666		666
667	# 尝试编译，预期抛出 ValueError 或相关配置错误	667	# 尝试编译，预期抛出 ValueError 或相关配置错误
668	# 注意：有些错误可能在编译阶段 (compile) 触发，有些可能在首次运行 (call) 时触发	668	# 注意：有些错误可能在编译阶段 (compile) 触发，有些可能在首次运行 (call) 时触发
669	try:	669	try:
670	compiled_fn = torch.compile(	670	compiled_fn = torch.compile(
671	model_fn,	671	model_fn,
672	backend='inductor',	672	backend='inductor',
673	dynamic=False,	673	dynamic=False,
674	options=invalid_shape_options	674	options=invalid_shape_options
675	)	675	)
676		676
677	# 准备输入数据触发编译	677	# 准备输入数据触发编译
678	A = torch.randn((2, 32), device=device)	678	A = torch.randn((2, 32), device=device)
679	B = torch.randn((2, 32), device=device)	679	B = torch.randn((2, 32), device=device)
680		680
681	with self.assertRaises((ValueError, RuntimeError, TypeError)) as cm:	681	with self.assertRaises((ValueError, RuntimeError, TypeError)) as cm:
682	compiled_fn(A, B)	682	compiled_fn(A, B)
683		683
684	print(f"\n[成功捕获异常]: {cm.exception}")	684	print(f"\n[成功捕获异常]: {cm.exception}")
685		685
686	except Exception as e:	686	except Exception as e:
687	# 如果在 torch.compile 阶段就直接崩溃，也算捕获成功	687	# 如果在 torch.compile 阶段就直接崩溃，也算捕获成功
688	print(f"\n[编译阶段直接报错]: {e}")	688	print(f"\n[编译阶段直接报错]: {e}")
689	self.assertIsInstance(e, (ValueError, RuntimeError, TypeError))	689	self.assertIsInstance(e, (ValueError, RuntimeError, TypeError))
690		690
691	def test_npu_malformed_option_type(self):	691	def test_npu_malformed_option_type(self):
692		692
693	# 构造格式完全错误的 options	693	# 构造格式完全错误的 options
694	malformed_options = {	694	malformed_options = {
695	"shape_handling_configs": "this should be a list, not a string"	695	"shape_handling_configs": "this should be a list, not a string"
696	}	696	}
697		697
698	with self.assertRaises(Exception):	698	with self.assertRaises(Exception):
699	compiled_fn = torch.compile(model_fn, options=malformed_options)	699	compiled_fn = torch.compile(model_fn, options=malformed_options)
700	compiled_fn(torch.randn(2, device="npu"), torch.randn(2, device="npu"))	700	compiled_fn(torch.randn(2, device="npu"), torch.randn(2, device="npu"))
701		701
702	def test_npu_shape_options_with_handling_disabled(self):	702	def test_npu_shape_options_with_handling_disabled(self):
703	# enable_shape_handling 设置为False, 同时保留主流的配置选项，测试是否会存在选项异常报错	703	# enable_shape_handling 设置为False, 同时保留主流的配置选项，测试是否会存在选项异常报错
704		704
705		705
706	shape_options_off = {	706	shape_options_off = {
707	"enable_shape_handling": False,	707	"enable_shape_handling": False,
708	"shape_handling_configs": [	708	"shape_handling_configs": [
709	{	709	{
710	"type": "BATCHSIZE",	710	"type": "BATCHSIZE",
711	"min_size": 1,	711	"min_size": 1,
712	"max_size": 1024,	712	"max_size": 1024,
713	"policy": "TIMES",	713	"policy": "TIMES",
714	}	714	}
715	]	715	]
716	}	716	}
717		717
718	# 尝试编译，预期可正确处理此类配置，则通过测试	718	# 尝试编译，预期可正确处理此类配置，则通过测试
719	try:	719	try:
720	compiled_fn = torch.compile(	720	compiled_fn = torch.compile(
721	model_fn,	721	model_fn,
722	backend='inductor',	722	backend='inductor',
723	dynamic=False,	723	dynamic=False,
724	options=shape_options_off	724	options=shape_options_off
725	)	725	)
726		726
727	# 准备输入数据触发编译	727	# 准备输入数据触发编译
728	A = torch.randn((2, 32), device=device)	728	A = torch.randn((2, 32), device=device)
729	B = torch.randn((2, 32), device=device)	729	B = torch.randn((2, 32), device=device)
730	compiled_fn(A, B)	730	compiled_fn(A, B)
731		731
732	except Exception as e:	732	except Exception as e:
733	self.fail(f"torch.compile raised {type(e).__name__} unexpectedly: {e}")	733	self.fail(f"torch.compile raised {type(e).__name__} unexpectedly: {e}")
734		734
735	def test_npu_shape_handling_whit_mutil_compile(self):	735	def test_npu_shape_handling_whit_mutil_compile(self):
736	# 多次编译同一个函数，验证 shape handling 是否正常工作	736	# 多次编译同一个函数，验证 shape handling 是否正常工作
737	# 1. 首次编译触发 shape handling	737	# 1. 首次编译触发 shape handling
738	# 2. 后续编译复用已有的 shape handling 逻辑	738	# 2. 后续编译复用已有的 shape handling 逻辑
739	try:	739	try:
740	compiled_fn = torch.compile(	740	compiled_fn = torch.compile(
741	model_fn,	741	model_fn,
742	backend='inductor',	742	backend='inductor',
743	dynamic=False,	743	dynamic=False,
744	options=shape_options	744	options=shape_options
745	)	745	)
746	# 第一次spilt	746	# 第一次spilt
747	compiled_fn(torch.randn((1025, 32), device=device), torch.randn((1025, 32), device=device))	747	compiled_fn(torch.randn((1025, 32), device=device), torch.randn((1025, 32), device=device))
748		748
749	# 第二次spilt	749	# 第二次spilt
750	compiled_fn(torch.randn((2048, 32), device=device), torch.randn((2048, 32), device=device))	750	compiled_fn(torch.randn((2048, 32), device=device), torch.randn((2048, 32), device=device))
751		751
752	# 第三次spilt	752	# 第三次spilt
753	compiled_fn(torch.randn((10240, 32), device=device), torch.randn((10240, 32), device=device))	753	compiled_fn(torch.randn((10240, 32), device=device), torch.randn((10240, 32), device=device))
754	except Exception as e:	754	except Exception as e:
755	self.fail(f"torch.compile raised {type(e).__name__} unexpectedly: ")	755	self.fail(f"torch.compile raised {type(e).__name__} unexpectedly: ")
756		756
757		757
758	class TestUnifiedCopy(TestCase):	758	class TestUnifiedCopy(TestCase):
759	def test_none_and_simple_types(self):	759	def test_none_and_simple_types(self):
760	"""Test unified_copy with None, list, dict, and tuple."""	760	"""Test unified_copy with None, list, dict, and tuple."""
761	# None	761	# None
762	self.assertIsNone(unified_copy(None))	762	self.assertIsNone(unified_copy(None))
763		763
764	# list	764	# list
765	lst = [1, 2, 3]	765	lst = [1, 2, 3]
766	copied_lst = unified_copy(lst)	766	copied_lst = unified_copy(lst)
767	self.assertEqual(lst, copied_lst)	767	self.assertEqual(lst, copied_lst)
768	self.assertIsNot(lst, copied_lst)	768	self.assertIsNot(lst, copied_lst)
769		769
770	# dict	770	# dict
771	d = {"a": 1, "b": [2, 3]}	771	d = {"a": 1, "b": [2, 3]}
772	copied_d = unified_copy(d)	772	copied_d = unified_copy(d)
773	self.assertEqual(d, copied_d)	773	self.assertEqual(d, copied_d)
774	self.assertIsNot(d, copied_d)	774	self.assertIsNot(d, copied_d)
775		775
776	# tuple	776	# tuple
777	t = (1, [2, 3])	777	t = (1, [2, 3])
778	copied_t = unified_copy(t)	778	copied_t = unified_copy(t)
779	self.assertEqual(t, copied_t)	779	self.assertEqual(t, copied_t)
780	self.assertIsInstance(copied_t, tuple)	780	self.assertIsInstance(copied_t, tuple)
781		781
782	def test_nested_structure_with_tensor(self):	782	def test_nested_structure_with_tensor(self):
783	"""Test nested structure containing NPU tensor."""	783	"""Test nested structure containing NPU tensor."""
784	original = {	784	original = {
785	"data": [torch.tensor([1.0, 2.0]), 42],	785	"data": [torch.tensor([1.0, 2.0]), 42],
786	"meta": {"shape": (2,)}	786	"meta": {"shape": (2,)}
787	}	787	}
788	copied = unified_copy(original)	788	copied = unified_copy(original)
789		789
790	# 结构和值一致	790	# 结构和值一致
791	self.assertEqual(original["meta"], copied["meta"])	791	self.assertEqual(original["meta"], copied["meta"])
792	self.assertTrue(torch.equal(original["data"][0], copied["data"][0]))	792	self.assertTrue(torch.equal(original["data"][0], copied["data"][0]))
793	self.assertEqual(original["data"][1], copied["data"][1])	793	self.assertEqual(original["data"][1], copied["data"][1])
794		794
795	# 验证独立副本	795	# 验证独立副本
796	original["data"][0][0] = 888.0	796	original["data"][0][0] = 888.0
797	self.assertNotEqual(copied["data"][0][0].item(), 888.0)	797	self.assertNotEqual(copied["data"][0][0].item(), 888.0)
798		798
799		799
800	def test_deepcopy_failure_returns_original(self):	800	def test_deepcopy_failure_returns_original(self):
801	class NonCopyable:	801	class NonCopyable:
802	def __deepcopy__(self, memo):	802	def __deepcopy__(self, memo):
803	raise TypeError("deepcopy not supported")	803	raise TypeError("deepcopy not supported")
804		804
805	obj = NonCopyable()	805	obj = NonCopyable()
806	copied = unified_copy(obj)	806	copied = unified_copy(obj)
807	self.assertIs(copied, obj)	807	self.assertIs(copied, obj)
808		808
809		809
810	instantiate_parametrized_tests(TestShapeHandling)	810	instantiate_parametrized_tests(TestShapeHandling)
811	instantiate_parametrized_tests(TestShapeHandlingBranchCoverage)	811	instantiate_parametrized_tests(TestShapeHandlingBranchCoverage)
812	instantiate_parametrized_tests(TestUnifiedCopy)	812	instantiate_parametrized_tests(TestUnifiedCopy)
813	instantiate_parametrized_tests(TestDynamicShapeCompile)	813	instantiate_parametrized_tests(TestDynamicShapeCompile)
814		814
815	if __name__ == '__main__':	815	if __name__ == '__main__':
816	run_tests()	816	run_tests()

Mtest/_inductor/test_tile_generator_32b_alignment.py+96-96

Mtest/_inductor/test_upcast_codegen.py+34-34

Mtest/_inductor/test_use_static_kernel.py+49-49

	@@ -1,50 +1,50 @@
1	import unittest	1	import unittest
2	import torch	2	import torch
3	import torch_npu	3	import torch_npu
4		4
5	from torch.testing._internal.common_utils import (	5	from torch.testing._internal.common_utils import (
6	run_tests,	6	run_tests,
7	parametrize,	7	parametrize,
8	instantiate_parametrized_tests,	8	instantiate_parametrized_tests,
9	)	9	)
10	from testutils import TestUtils	10	from testutils import TestUtils
11		11
12		12
13	class TestInductorStaticKernel(TestUtils):	13	class TestInductorStaticKernel(TestUtils):
14		14
15	def simple_op(self, x):	15	def simple_op(self, x):
16	return torch.neg(x)	16	return torch.neg(x)
17		17
18	@parametrize("shape", [(1024, 1024), (4096,)])	18	@parametrize("shape", [(1024, 1024), (4096,)])
19	@parametrize("dtype", [torch.float16, torch.float32])	19	@parametrize("dtype", [torch.float16, torch.float32])
20	def test_inductor_static_kernel(self, shape, dtype):	20	def test_inductor_static_kernel(self, shape, dtype):
21	device = "npu"	21	device = "npu"
22		22
23	x = torch.randn(shape, dtype=dtype, device=device)	23	x = torch.randn(shape, dtype=dtype, device=device)
24		24
25	ref = self.simple_op(x)	25	ref = self.simple_op(x)
26		26
27	torch._inductor.config.triton.cudagraph_trees = False	27	torch._inductor.config.triton.cudagraph_trees = False
28	torch_npu.npu.aclnn._use_static_aclnn_kernel = True	28	torch_npu.npu.aclnn._use_static_aclnn_kernel = True
29		29
30	compiled_fn = torch.compile(	30	compiled_fn = torch.compile(
31	self.simple_op,	31	self.simple_op,
32	backend="inductor",	32	backend="inductor",
33	dynamic=False	33	dynamic=False
34	)	34	)
35		35
36	for _ in range(3):	36	for _ in range(3):
37	compiled_fn(x)	37	compiled_fn(x)
38		38
39	torch.npu.synchronize()	39	torch.npu.synchronize()
40		40
41	out = compiled_fn(x)	41	out = compiled_fn(x)
42	torch.npu.synchronize()	42	torch.npu.synchronize()
43		43
44	self.assertEqual(ref, out)	44	self.assertEqual(ref, out)
45		45
46	instantiate_parametrized_tests(TestInductorStaticKernel)	46	instantiate_parametrized_tests(TestInductorStaticKernel)
47		47
48	if __name__ == "__main__":	48	if __name__ == "__main__":
49	torch.npu.config.allow_internal_format = False	49	torch.npu.config.allow_internal_format = False
50	run_tests()	50	run_tests()

Mtest/_inductor/testutils.py+0-1

	@@ -73,4 +73,3 @@ class BenchmarkTestUtils(TestCase):
73	writer.writeheader()	73	writer.writeheader()
74		74
75	writer.writerow(perf_info)	75	writer.writerow(perf_info)
76

Mtest/autograd/test_autograd_fallback.py+30-30

	@@ -1,30 +1,30 @@
1	import torch	1	import torch
2	from torch.testing._internal.common_utils import (	2	from torch.testing._internal.common_utils import (
3	run_tests,	3	run_tests,
4	TestCase,	4	TestCase,
5	)	5	)
6	import torch_npu	6	import torch_npu
7		7
8	class TestAutogradFallback(TestCase):	8	class TestAutogradFallback(TestCase):
9		9
10	def test_pad_backward_warn(self):	10	def test_pad_backward_warn(self):
11		11
12	def _exec_npu_pad():	12	def _exec_npu_pad():
13	npu_input = torch.randn(2, 3).npu()	13	npu_input = torch.randn(2, 3).npu()
14	npu_input.requires_grad = True	14	npu_input.requires_grad = True
15	pads = (1, 1, 1, 1)	15	pads = (1, 1, 1, 1)
16	output = torch_npu.npu_pad(npu_input, pads)	16	output = torch_npu.npu_pad(npu_input, pads)
17	output.backward(torch.ones_like(output))	17	output.backward(torch.ones_like(output))
18		18
19	# When set to "nothing," calling the reverse function directly causes an error.	19	# When set to "nothing," calling the reverse function directly causes an error.
20	torch._C._set_autograd_fallback_mode("nothing")	20	torch._C._set_autograd_fallback_mode("nothing")
21	with self.assertRaisesRegex(RuntimeError, "does not require grad"):	21	with self.assertRaisesRegex(RuntimeError, "does not require grad"):
22	_exec_npu_pad()	22	_exec_npu_pad()
23		23
24	# When set to "warn," calling the print function emits a warning.	24	# When set to "warn," calling the print function emits a warning.
25	torch._C._set_autograd_fallback_mode("warn")	25	torch._C._set_autograd_fallback_mode("warn")
26	_exec_npu_pad()	26	_exec_npu_pad()
27		27
28		28
29	if __name__ == "__main__":	29	if __name__ == "__main__":
30	run_tests()	30	run_tests()

Mtest/custom_ops/test_fast_gelu.py+41-41

Mtest/custom_ops/test_npu_confusion_transpose.py+35-35

	@@ -1,35 +1,35 @@
1	import numpy as np	1	import numpy as np
2	import torch	2	import torch
3		3
4	import torch_npu	4	import torch_npu
5	from torch_npu.testing.testcase import TestCase, run_tests	5	from torch_npu.testing.testcase import TestCase, run_tests
6	from torch_npu.testing.common_utils import create_common_tensor	6	from torch_npu.testing.common_utils import create_common_tensor
7		7
8		8
9	class TestConfusionTranspose(TestCase):	9	class TestConfusionTranspose(TestCase):
10		10
11	def supported_op_exec(self, input1, perm, shape, transpose_first):	11	def supported_op_exec(self, input1, perm, shape, transpose_first):
12	if transpose_first:	12	if transpose_first:
13	output = input1.permute(*perm).contiguous().view(shape)	13	output = input1.permute(*perm).contiguous().view(shape)
14	else:	14	else:
15	output = input1.view(shape).permute(*perm)	15	output = input1.view(shape).permute(*perm)
16	return output.cpu().detach()	16	return output.cpu().detach()
17		17
18	def custom_op_exec(self, input1, perm, shape, transpose_first):	18	def custom_op_exec(self, input1, perm, shape, transpose_first):
19	output = torch_npu.npu_confusion_transpose(input1, perm, shape, transpose_first)	19	output = torch_npu.npu_confusion_transpose(input1, perm, shape, transpose_first)
20	return output.cpu().detach()	20	return output.cpu().detach()
21		21
22	def test_npu_confusion_transpose(self, device="npu"):	22	def test_npu_confusion_transpose(self, device="npu"):
23	item = [np.float32, 0, [1, 576, 2560]]	23	item = [np.float32, 0, [1, 576, 2560]]
24	_, npu_input = create_common_tensor(item, 0, 100)	24	_, npu_input = create_common_tensor(item, 0, 100)
25	perm = (0, 2, 1, 3)	25	perm = (0, 2, 1, 3)
26	shape = [1, 576, 32, 80]	26	shape = [1, 576, 32, 80]
27	transpose_first = False	27	transpose_first = False
28		28
29	supported_output = self.supported_op_exec(npu_input, perm, shape, transpose_first)	29	supported_output = self.supported_op_exec(npu_input, perm, shape, transpose_first)
30	custom_output = self.custom_op_exec(npu_input, perm, shape, transpose_first)	30	custom_output = self.custom_op_exec(npu_input, perm, shape, transpose_first)
31	self.assertRtolEqual(supported_output, custom_output)	31	self.assertRtolEqual(supported_output, custom_output)
32		32
33		33
34	if __name__ == "__main__":	34	if __name__ == "__main__":
35	run_tests()	35	run_tests()

Mtest/custom_ops/test_npu_convolution.py+52-52

	@@ -1,52 +1,52 @@
1	#	1	#
2		2
3	import numpy as np	3	import numpy as np
4	import torch	4	import torch
5		5
6	import torch_npu	6	import torch_npu
7	from torch_npu.testing.testcase import TestCase, run_tests	7	from torch_npu.testing.testcase import TestCase, run_tests
8	from torch_npu.testing.common_utils import create_common_tensor	8	from torch_npu.testing.common_utils import create_common_tensor
9		9
10		10
11	class TestConvolution(TestCase):	11	class TestConvolution(TestCase):
12		12
13	def supported_op_exec(self, input1, weight, bias, stride, padding, dilation, groups):	13	def supported_op_exec(self, input1, weight, bias, stride, padding, dilation, groups):
14	dim = input1.dim()	14	dim = input1.dim()
15	if dim == 4:	15	if dim == 4:
16	output = torch.nn.functional.conv2d(input1, weight, bias, stride, padding, dilation, groups)	16	output = torch.nn.functional.conv2d(input1, weight, bias, stride, padding, dilation, groups)
17	if dim == 5:	17	if dim == 5:
18	is_dilated = False	18	is_dilated = False
19	for d in dilation:	19	for d in dilation:
20	is_dilated \|= (d != 1)	20	is_dilated \|= (d != 1)
21	if groups == 1 and not is_dilated:	21	if groups == 1 and not is_dilated:
22	kernel_size = weight.size()[2]	22	kernel_size = weight.size()[2]
23	output = torch._C._nn.slow_conv3d(input1, weight, kernel_size, bias, stride, padding)	23	output = torch._C._nn.slow_conv3d(input1, weight, kernel_size, bias, stride, padding)
24	else:	24	else:
25	output = torch.nn.functional.conv3d(input1, weight, bias, stride, padding, dilation, groups)	25	output = torch.nn.functional.conv3d(input1, weight, bias, stride, padding, dilation, groups)
26	return output.cpu().detach()	26	return output.cpu().detach()
27		27
28	def custom_op_exec(self, input1, weight, bias, stride, padding, dilation, groups):	28	def custom_op_exec(self, input1, weight, bias, stride, padding, dilation, groups):
29	output = torch_npu.npu_convolution(input1, weight, bias, stride, padding, dilation, groups)	29	output = torch_npu.npu_convolution(input1, weight, bias, stride, padding, dilation, groups)
30	return output.cpu().detach()	30	return output.cpu().detach()
31		31
32	def test_npu_convolution(self, device="npu"):	32	def test_npu_convolution(self, device="npu"):
33	items = [[[np.float32, 0, [16, 128, 112, 112]], [np.float32, 0, [256, 128, 3, 3]], [np.float32, 2, [256]],	33	items = [[[np.float32, 0, [16, 128, 112, 112]], [np.float32, 0, [256, 128, 3, 3]], [np.float32, 2, [256]],
34	[1, 1], [1, 1], [1, 1], 1],	34	[1, 1], [1, 1], [1, 1], 1],
35	[[np.float16, 30, [1, 128, 4, 14, 14]], [np.float16, 30, [1, 128, 3, 3, 3]], None,	35	[[np.float16, 30, [1, 128, 4, 14, 14]], [np.float16, 30, [1, 128, 3, 3, 3]], None,
36	[1, 1, 1], [1, 1, 1], [1, 1, 1], 1]]	36	[1, 1, 1], [1, 1, 1], [1, 1, 1], 1]]
37	for item in items:	37	for item in items:
38	_, npu_input = create_common_tensor(item[0], -1, 1)	38	_, npu_input = create_common_tensor(item[0], -1, 1)
39	_, weight = create_common_tensor(item[1], -1, 1)	39	_, weight = create_common_tensor(item[1], -1, 1)
40	_, bias = create_common_tensor(item[2], -1, 1) if item[2] else _, None	40	_, bias = create_common_tensor(item[2], -1, 1) if item[2] else _, None
41	stride = item[3]	41	stride = item[3]
42	padding = item[4]	42	padding = item[4]
43	dilation = item[5]	43	dilation = item[5]
44	groups = item[6]	44	groups = item[6]
45		45
46	supported_output = self.supported_op_exec(npu_input, weight, bias, stride, padding, dilation, groups)	46	supported_output = self.supported_op_exec(npu_input, weight, bias, stride, padding, dilation, groups)
47	custom_output = self.custom_op_exec(npu_input, weight, bias, stride, padding, dilation, groups)	47	custom_output = self.custom_op_exec(npu_input, weight, bias, stride, padding, dilation, groups)
48	self.assertRtolEqual(supported_output, custom_output)	48	self.assertRtolEqual(supported_output, custom_output)
49		49
50		50
51	if __name__ == "__main__":	51	if __name__ == "__main__":
52	run_tests()	52	run_tests()

Mtest/custom_ops/test_npu_dtype_cast.py+50-50

Mtest/custom_ops/test_npu_format_cast.py+31-31

	@@ -1,31 +1,31 @@
1	import numpy as np	1	import numpy as np
2	import torch	2	import torch
3		3
4	import torch_npu	4	import torch_npu
5	from torch_npu.testing.testcase import TestCase, run_tests	5	from torch_npu.testing.testcase import TestCase, run_tests
6	from torch_npu.testing.common_utils import create_common_tensor	6	from torch_npu.testing.common_utils import create_common_tensor
7		7
8		8
9	class TestFormatCast(TestCase):	9	class TestFormatCast(TestCase):
10		10
11	def supported_op_exec(self, input1):	11	def supported_op_exec(self, input1):
12	m = torch.nn.Identity(54, unused_argument1=0.1, unused_argument2=False)	12	m = torch.nn.Identity(54, unused_argument1=0.1, unused_argument2=False)
13	output = m(input1)	13	output = m(input1)
14	return output.cpu().detach()	14	return output.cpu().detach()
15		15
16	def custom_op_exec(self, input1, acl_format):	16	def custom_op_exec(self, input1, acl_format):
17	output = torch_npu.npu_format_cast(input1, acl_format)	17	output = torch_npu.npu_format_cast(input1, acl_format)
18	return output.cpu().detach()	18	return output.cpu().detach()
19		19
20	def test_npu_format_cast(self, device="npu"):	20	def test_npu_format_cast(self, device="npu"):
21	item = [np.float16, 0, (2, 2, 4, 4)]	21	item = [np.float16, 0, (2, 2, 4, 4)]
22	_, npu_input = create_common_tensor(item, -1, 1)	22	_, npu_input = create_common_tensor(item, -1, 1)
23	acl_format = 3	23	acl_format = 3
24		24
25	supported_output = self.supported_op_exec(npu_input)	25	supported_output = self.supported_op_exec(npu_input)
26	custom_output = self.custom_op_exec(npu_input, acl_format)	26	custom_output = self.custom_op_exec(npu_input, acl_format)
27	self.assertRtolEqual(supported_output, custom_output)	27	self.assertRtolEqual(supported_output, custom_output)
28		28
29		29
30	if __name__ == "__main__":	30	if __name__ == "__main__":
31	run_tests()	31	run_tests()

Mtest/custom_ops/test_npu_grid_assign_positive.py+44-44

Mtest/custom_ops/test_npu_ps_roi_pooling.py+94-94

Mtest/custom_ops/test_npu_softmax_cross_entropy_with_logits.py+31-31

Mtest/custom_ops/test_scatter_update.py+115-115

	@@ -1,115 +1,115 @@
1	import numpy as np	1	import numpy as np
2	import torch	2	import torch
3		3
4	import torch_npu	4	import torch_npu
5	from torch_npu.testing.testcase import TestCase, run_tests	5	from torch_npu.testing.testcase import TestCase, run_tests
6	from torch_npu.testing.common_utils import create_common_tensor	6	from torch_npu.testing.common_utils import create_common_tensor
7		7
8		8
9	class TestScatterUpdate(TestCase):	9	class TestScatterUpdate(TestCase):
10		10
11	def supported_scatter_update_exec(self, var, updates, start, length, axis=-2):	11	def supported_scatter_update_exec(self, var, updates, start, length, axis=-2):
12	var_input = var.clone()	12	var_input = var.clone()
13	var_input[:, :, start: start + length, :] = updates	13	var_input[:, :, start: start + length, :] = updates
14	return var_input	14	return var_input
15		15
16	def custom_scatter_update_exec(self, var, indices, updates, axis=-2):	16	def custom_scatter_update_exec(self, var, indices, updates, axis=-2):
17	output = torch_npu.scatter_update(var, indices, updates, axis)	17	output = torch_npu.scatter_update(var, indices, updates, axis)
18	return output.cpu().detach()	18	return output.cpu().detach()
19		19
20		20
21	def supported_scatter_update__exec(self, var, updates, start, length, axis=-2):	21	def supported_scatter_update__exec(self, var, updates, start, length, axis=-2):
22	var[:, :, start: start + length, :] = updates	22	var[:, :, start: start + length, :] = updates
23	return var	23	return var
24		24
25	def custom_scatter_update__exec(self, var, indices, updates, axis=-2):	25	def custom_scatter_update__exec(self, var, indices, updates, axis=-2):
26	torch_npu.scatter_update_(var, indices, updates, axis)	26	torch_npu.scatter_update_(var, indices, updates, axis)
27	return var.cpu().detach()	27	return var.cpu().detach()
28		28
29	def custom_scatter_update__exec_return(self, var, indices, updates, axis=-2):	29	def custom_scatter_update__exec_return(self, var, indices, updates, axis=-2):
30	result = torch_npu.scatter_update_(var, indices, updates, axis)	30	result = torch_npu.scatter_update_(var, indices, updates, axis)
31	return result.cpu().detach()	31	return result.cpu().detach()
32		32
33		33
34	def test_scatter_update(self, device="npu"):	34	def test_scatter_update(self, device="npu"):
35	# input_dtype, slice_start, slice_length	35	# input_dtype, slice_start, slice_length
36	items = [	36	items = [
37	[torch.float16, 0, 1],	37	[torch.float16, 0, 1],
38	[torch.float16, 2, 1],	38	[torch.float16, 2, 1],
39	[torch.float16, 4, 8],	39	[torch.float16, 4, 8],
40	[torch.float16, 0, 16],	40	[torch.float16, 0, 16],
41	[torch.float32, 0, 1],	41	[torch.float32, 0, 1],
42	[torch.float32, 2, 1],	42	[torch.float32, 2, 1],
43	[torch.float32, 4, 8],	43	[torch.float32, 4, 8],
44	[torch.float32, 0, 16],	44	[torch.float32, 0, 16],
45	]	45	]
46		46
47	for item in items:	47	for item in items:
48	in_self_cpu = torch.randn(4, 8, 16, 64, dtype=item[0])	48	in_self_cpu = torch.randn(4, 8, 16, 64, dtype=item[0])
49	in_self_npu = in_self_cpu.npu()	49	in_self_npu = in_self_cpu.npu()
50	in_update_cpu = torch.randn(4, 8, item[2], 64, dtype=item[0])	50	in_update_cpu = torch.randn(4, 8, item[2], 64, dtype=item[0])
51	in_update_npu = in_update_cpu.npu()	51	in_update_npu = in_update_cpu.npu()
52	in_indices_npu = torch.tensor([item[1], item[1], item[1], item[1]]).npu()	52	in_indices_npu = torch.tensor([item[1], item[1], item[1], item[1]]).npu()
53		53
54	supported_output = self.supported_scatter_update_exec(in_self_cpu, in_update_cpu, item[1], item[2])	54	supported_output = self.supported_scatter_update_exec(in_self_cpu, in_update_cpu, item[1], item[2])
55	custom_output = self.custom_scatter_update_exec(in_self_npu, in_indices_npu, in_update_npu)	55	custom_output = self.custom_scatter_update_exec(in_self_npu, in_indices_npu, in_update_npu)
56	self.assertRtolEqual(supported_output, custom_output)	56	self.assertRtolEqual(supported_output, custom_output)
57	# check whether the custom operator modifies the input_self	57	# check whether the custom operator modifies the input_self
58	self.assertRtolEqual(in_self_cpu, in_self_npu.cpu().detach())	58	self.assertRtolEqual(in_self_cpu, in_self_npu.cpu().detach())
59		59
60		60
61	def test_scatter_update_(self, device="npu"):	61	def test_scatter_update_(self, device="npu"):
62	# input_dtype, slice_start, slice_length	62	# input_dtype, slice_start, slice_length
63	items = [	63	items = [
64	[torch.float16, 0, 1],	64	[torch.float16, 0, 1],
65	[torch.float16, 4, 1],	65	[torch.float16, 4, 1],
66	[torch.float16, 8, 32],	66	[torch.float16, 8, 32],
67	[torch.float16, 0, 64],	67	[torch.float16, 0, 64],
68	[torch.float32, 0, 1],	68	[torch.float32, 0, 1],
69	[torch.float32, 4, 1],	69	[torch.float32, 4, 1],
70	[torch.float32, 8, 32],	70	[torch.float32, 8, 32],
71	[torch.float32, 0, 64],	71	[torch.float32, 0, 64],
72	]	72	]
73	for item in items:	73	for item in items:
74	in_self_cpu = torch.randn(4, 8, 64, 128, dtype=item[0])	74	in_self_cpu = torch.randn(4, 8, 64, 128, dtype=item[0])
75	in_self_npu = in_self_cpu.npu()	75	in_self_npu = in_self_cpu.npu()
76	in_update_cpu = torch.randn(4, 8, item[2], 128, dtype=item[0])	76	in_update_cpu = torch.randn(4, 8, item[2], 128, dtype=item[0])
77	in_update_npu = in_update_cpu.npu()	77	in_update_npu = in_update_cpu.npu()
78	in_indices_npu = torch.tensor([item[1], item[1], item[1], item[1]]).npu()	78	in_indices_npu = torch.tensor([item[1], item[1], item[1], item[1]]).npu()
79		79
80	supported_output = self.supported_scatter_update__exec(in_self_cpu, in_update_cpu, item[1], item[2])	80	supported_output = self.supported_scatter_update__exec(in_self_cpu, in_update_cpu, item[1], item[2])
81	custom_output = self.custom_scatter_update__exec(in_self_npu, in_indices_npu, in_update_npu)	81	custom_output = self.custom_scatter_update__exec(in_self_npu, in_indices_npu, in_update_npu)
82	self.assertRtolEqual(supported_output, custom_output)	82	self.assertRtolEqual(supported_output, custom_output)
83	# check whether the custom operator modifies the input_self	83	# check whether the custom operator modifies the input_self
84	self.assertRtolEqual(in_self_cpu, in_self_npu.cpu().detach())	84	self.assertRtolEqual(in_self_cpu, in_self_npu.cpu().detach())
85		85
86		86
87	def test_scatter_update__return(self, device="npu"):	87	def test_scatter_update__return(self, device="npu"):
88	# input_dtype, slice_start, slice_length	88	# input_dtype, slice_start, slice_length
89	items = [	89	items = [
90	[torch.float16, 0, 1],	90	[torch.float16, 0, 1],
91	[torch.float16, 8, 1],	91	[torch.float16, 8, 1],
92	[torch.float16, 16, 8],	92	[torch.float16, 16, 8],
93	[torch.float16, 0, 32],	93	[torch.float16, 0, 32],
94	[torch.float32, 0, 1],	94	[torch.float32, 0, 1],
95	[torch.float32, 8, 1],	95	[torch.float32, 8, 1],
96	[torch.float32, 16, 8],	96	[torch.float32, 16, 8],
97	[torch.float32, 0, 32],	97	[torch.float32, 0, 32],
98	]	98	]
99	for item in items:	99	for item in items:
100	in_self_cpu = torch.randn(12, 2, 64, 128, dtype=item[0])	100	in_self_cpu = torch.randn(12, 2, 64, 128, dtype=item[0])
101	in_self_npu = in_self_cpu.npu()	101	in_self_npu = in_self_cpu.npu()
102	in_update_cpu = torch.randn(12, 2, item[2], 128, dtype=item[0])	102	in_update_cpu = torch.randn(12, 2, item[2], 128, dtype=item[0])
103	in_update_npu = in_update_cpu.npu()	103	in_update_npu = in_update_cpu.npu()
104	in_indices_npu = torch.tensor([item[1], item[1], item[1], item[1], item[1], item[1],	104	in_indices_npu = torch.tensor([item[1], item[1], item[1], item[1], item[1], item[1],
105	item[1], item[1], item[1], item[1], item[1], item[1]]).npu()	105	item[1], item[1], item[1], item[1], item[1], item[1]]).npu()
106		106
107	supported_output = self.supported_scatter_update__exec(in_self_cpu, in_update_cpu, item[1], item[2])	107	supported_output = self.supported_scatter_update__exec(in_self_cpu, in_update_cpu, item[1], item[2])
108	custom_output = self.custom_scatter_update__exec(in_self_npu, in_indices_npu, in_update_npu)	108	custom_output = self.custom_scatter_update__exec(in_self_npu, in_indices_npu, in_update_npu)
109	self.assertRtolEqual(supported_output, custom_output)	109	self.assertRtolEqual(supported_output, custom_output)
110	# check whether the custom operator modifies the input_self	110	# check whether the custom operator modifies the input_self
111	self.assertRtolEqual(in_self_cpu, in_self_npu.cpu().detach())	111	self.assertRtolEqual(in_self_cpu, in_self_npu.cpu().detach())
112		112
113		113
114	if __name__ == "__main__":	114	if __name__ == "__main__":
115	run_tests()	115	run_tests()

Mtest/distributed/elastic/events/test_events_api.py+319-319

	@@ -1,320 +1,320 @@
1	"""	1	"""
2	Add validation cases for torch.distributed.elastic.events.record API:	2	Add validation cases for torch.distributed.elastic.events.record API:
3		3
4	1. PyTorch community lacks direct test cases for torch.distributed.elastic.events.record	4	1. PyTorch community lacks direct test cases for torch.distributed.elastic.events.record
5	in the standard test suite, so this file is added.	5	in the standard test suite, so this file is added.
6		6
7	2. This file validates the following APIs:	7	2. This file validates the following APIs:
8	torch.distributed.elastic.events.record (This is a pure Python-level event logging API with no hardware dependency.)	8	torch.distributed.elastic.events.record (This is a pure Python-level event logging API with no hardware dependency.)
9	torch.distributed.elastic.events.api.EventMetadataValue (Type alias for metadata values)	9	torch.distributed.elastic.events.api.EventMetadataValue (Type alias for metadata values)
10	(extendable)	10	(extendable)
11	"""	11	"""
12		12
13	import json	13	import json
14	import time	14	import time
15	from typing import Union, Optional, get_args, get_origin	15	from typing import Union, Optional, get_args, get_origin
16	from unittest.mock import patch, MagicMock	16	from unittest.mock import patch, MagicMock
17		17
18	import torch	18	import torch
19	from torch.distributed.elastic.events import record, get_logging_handler	19	from torch.distributed.elastic.events import record, get_logging_handler
20	from torch.distributed.elastic.events.api import Event, EventSource, EventMetadataValue	20	from torch.distributed.elastic.events.api import Event, EventSource, EventMetadataValue
21	from torch.testing._internal.common_utils import TestCase, run_tests	21	from torch.testing._internal.common_utils import TestCase, run_tests
22		22
23		23
24	class TestEventsRecord(TestCase):	24	class TestEventsRecord(TestCase):
25	"""Test torch.distributed.elastic.events.record method."""	25	"""Test torch.distributed.elastic.events.record method."""
26		26
27	DESTINATION_NULL = "null"	27	DESTINATION_NULL = "null"
28	DESTINATION_CONSOLE = "console"	28	DESTINATION_CONSOLE = "console"
29	TEST_EVENT_PREFIX = "test_event_"	29	TEST_EVENT_PREFIX = "test_event_"
30		30
31	def tearDown(self):	31	def tearDown(self):
32	"""Clean up resources after each test case to ensure test isolation."""	32	"""Clean up resources after each test case to ensure test isolation."""
33	from torch.distributed.elastic.events import _events_loggers	33	from torch.distributed.elastic.events import _events_loggers
34	_events_loggers.clear()	34	_events_loggers.clear()
35		35
36	if hasattr(get_logging_handler, "cache_clear"):	36	if hasattr(get_logging_handler, "cache_clear"):
37	get_logging_handler.cache_clear()	37	get_logging_handler.cache_clear()
38		38
39	patch.stopall()	39	patch.stopall()
40	super().tearDown()	40	super().tearDown()
41		41
42	def test_record_null_destination(self):	42	def test_record_null_destination(self):
43	"""Verify that record does not raise an exception when using the default null destination."""	43	"""Verify that record does not raise an exception when using the default null destination."""
44	event = Event(	44	event = Event(
45	name=f"{self.TEST_EVENT_PREFIX}null",	45	name=f"{self.TEST_EVENT_PREFIX}null",
46	source=EventSource.WORKER,	46	source=EventSource.WORKER,
47	metadata={"key": "value"}	47	metadata={"key": "value"}
48	)	48	)
49	record(event)	49	record(event)
50		50
51	def test_record_console_destination(self):	51	def test_record_console_destination(self):
52	"""Verify that record does not raise an exception when using the console destination."""	52	"""Verify that record does not raise an exception when using the console destination."""
53	event = Event(	53	event = Event(
54	name=f"{self.TEST_EVENT_PREFIX}console",	54	name=f"{self.TEST_EVENT_PREFIX}console",
55	source=EventSource.AGENT,	55	source=EventSource.AGENT,
56	metadata={"stage": "init"}	56	metadata={"stage": "init"}
57	)	57	)
58	record(event, destination=self.DESTINATION_CONSOLE)	58	record(event, destination=self.DESTINATION_CONSOLE)
59		59
60	def test_record_with_timestamp(self):	60	def test_record_with_timestamp(self):
61	"""Verify that record can correctly log an event with a custom or auto-generated timestamp."""	61	"""Verify that record can correctly log an event with a custom or auto-generated timestamp."""
62	custom_timestamp = int(time.time() * 1000)	62	custom_timestamp = int(time.time() * 1000)
63	event = Event(	63	event = Event(
64	name=f"{self.TEST_EVENT_PREFIX}timestamp_custom",	64	name=f"{self.TEST_EVENT_PREFIX}timestamp_custom",
65	source=EventSource.WORKER,	65	source=EventSource.WORKER,
66	timestamp=custom_timestamp,	66	timestamp=custom_timestamp,
67	metadata={"ts": custom_timestamp}	67	metadata={"ts": custom_timestamp}
68	)	68	)
69	record(event, destination=self.DESTINATION_NULL)	69	record(event, destination=self.DESTINATION_NULL)
70	self.assertEqual(event.timestamp, custom_timestamp)	70	self.assertEqual(event.timestamp, custom_timestamp)
71	self.assertIsInstance(event.timestamp, int)	71	self.assertIsInstance(event.timestamp, int)
72	self.assertGreaterEqual(event.timestamp, 0)	72	self.assertGreaterEqual(event.timestamp, 0)
73		73
74	event_with_auto_timestamp = Event(	74	event_with_auto_timestamp = Event(
75	name=f"{self.TEST_EVENT_PREFIX}timestamp_auto",	75	name=f"{self.TEST_EVENT_PREFIX}timestamp_auto",
76	source=EventSource.WORKER	76	source=EventSource.WORKER
77	)	77	)
78	record(event_with_auto_timestamp)	78	record(event_with_auto_timestamp)
79	self.assertIsNotNone(event_with_auto_timestamp.timestamp)	79	self.assertIsNotNone(event_with_auto_timestamp.timestamp)
80	self.assertIsInstance(event_with_auto_timestamp.timestamp, int)	80	self.assertIsInstance(event_with_auto_timestamp.timestamp, int)
81	self.assertGreaterEqual(event_with_auto_timestamp.timestamp, 0)	81	self.assertGreaterEqual(event_with_auto_timestamp.timestamp, 0)
82		82
83	@patch("torch.distributed.elastic.events._get_or_create_logger")	83	@patch("torch.distributed.elastic.events._get_or_create_logger")
84	def test_record_with_various_metadata_types(self, mock_get_logger):	84	def test_record_with_various_metadata_types(self, mock_get_logger):
85	"""Verify that record correctly passes various metadata types to the underlying logger."""	85	"""Verify that record correctly passes various metadata types to the underlying logger."""
86	mock_logger = MagicMock()	86	mock_logger = MagicMock()
87	mock_get_logger.return_value = mock_logger	87	mock_get_logger.return_value = mock_logger
88		88
89	metadata = {	89	metadata = {
90	"str_val": "hello",	90	"str_val": "hello",
91	"int_val": 42,	91	"int_val": 42,
92	"float_val": 3.14,	92	"float_val": 3.14,
93	"bool_val": True,	93	"bool_val": True,
94	"none_val": None,	94	"none_val": None,
95	"list_val": [1, 2, 3],	95	"list_val": [1, 2, 3],
96	"nested_dict": {"sub_key": "sub_val"}	96	"nested_dict": {"sub_key": "sub_val"}
97	}	97	}
98	event = Event(	98	event = Event(
99	name=f"{self.TEST_EVENT_PREFIX}metadata_types",	99	name=f"{self.TEST_EVENT_PREFIX}metadata_types",
100	source=EventSource.AGENT,	100	source=EventSource.AGENT,
101	metadata=metadata	101	metadata=metadata
102	)	102	)
103	record(event, destination=self.DESTINATION_CONSOLE)	103	record(event, destination=self.DESTINATION_CONSOLE)
104		104
105	# Verify the internal call chain	105	# Verify the internal call chain
106	mock_get_logger.assert_called_once_with(self.DESTINATION_CONSOLE)	106	mock_get_logger.assert_called_once_with(self.DESTINATION_CONSOLE)
107	mock_logger.info.assert_called_once()	107	mock_logger.info.assert_called_once()
108		108
109	# Verify that the serialized metadata is correctly preserved	109	# Verify that the serialized metadata is correctly preserved
110	logged_raw = mock_logger.info.call_args[0][0]	110	logged_raw = mock_logger.info.call_args[0][0]
111	logged_data = json.loads(logged_raw)	111	logged_data = json.loads(logged_raw)
112	self.assertEqual(logged_data["metadata"], metadata)	112	self.assertEqual(logged_data["metadata"], metadata)
113		113
114	def test_record_multiple_events(self):	114	def test_record_multiple_events(self):
115	"""Verify that consecutive calls to record do not interfere with each other."""	115	"""Verify that consecutive calls to record do not interfere with each other."""
116	event1 = Event(	116	event1 = Event(
117	name=f"{self.TEST_EVENT_PREFIX}multiple_1",	117	name=f"{self.TEST_EVENT_PREFIX}multiple_1",
118	source=EventSource.WORKER,	118	source=EventSource.WORKER,
119	metadata={"seq": 1}	119	metadata={"seq": 1}
120	)	120	)
121	event2 = Event(	121	event2 = Event(
122	name=f"{self.TEST_EVENT_PREFIX}multiple_2",	122	name=f"{self.TEST_EVENT_PREFIX}multiple_2",
123	source=EventSource.WORKER,	123	source=EventSource.WORKER,
124	metadata={"seq": 2}	124	metadata={"seq": 2}
125	)	125	)
126		126
127	record(event1)	127	record(event1)
128	record(event2)	128	record(event2)
129		129
130	self.assertIn("seq", event1.metadata)	130	self.assertIn("seq", event1.metadata)
131	self.assertEqual(event1.metadata["seq"], 1)	131	self.assertEqual(event1.metadata["seq"], 1)
132	self.assertIn("seq", event2.metadata)	132	self.assertIn("seq", event2.metadata)
133	self.assertEqual(event2.metadata["seq"], 2)	133	self.assertEqual(event2.metadata["seq"], 2)
134		134
135	@patch("torch.distributed.elastic.events._get_or_create_logger")	135	@patch("torch.distributed.elastic.events._get_or_create_logger")
136	def test_record_calls_get_or_create_logger(self, mock_get_logger):	136	def test_record_calls_get_or_create_logger(self, mock_get_logger):
137	"""	137	"""
138	Verify that record internally calls _get_or_create_logger to obtain the logger,	138	Verify that record internally calls _get_or_create_logger to obtain the logger,
139	and then calls .info() on that logger with the serialized event.	139	and then calls .info() on that logger with the serialized event.
140	"""	140	"""
141	mock_logger = MagicMock()	141	mock_logger = MagicMock()
142	mock_get_logger.return_value = mock_logger	142	mock_get_logger.return_value = mock_logger
143		143
144	event = Event(	144	event = Event(
145	name=f"{self.TEST_EVENT_PREFIX}mock",	145	name=f"{self.TEST_EVENT_PREFIX}mock",
146	source=EventSource.WORKER	146	source=EventSource.WORKER
147	)	147	)
148	record(event, destination=self.DESTINATION_CONSOLE)	148	record(event, destination=self.DESTINATION_CONSOLE)
149		149
150	mock_get_logger.assert_called_once_with(self.DESTINATION_CONSOLE)	150	mock_get_logger.assert_called_once_with(self.DESTINATION_CONSOLE)
151	mock_logger.info.assert_called_once_with(event.serialize())	151	mock_logger.info.assert_called_once_with(event.serialize())
152		152
153	def test_record_event_name_empty(self):	153	def test_record_event_name_empty(self):
154	"""Verify that record does not raise an exception when the event name is an empty string."""	154	"""Verify that record does not raise an exception when the event name is an empty string."""
155	event = Event(name="", source=EventSource.WORKER)	155	event = Event(name="", source=EventSource.WORKER)
156	record(event)	156	record(event)
157		157
158	def test_record_event_name_special_chars(self):	158	def test_record_event_name_special_chars(self):
159	"""Verify event names with special characters (spaces, symbols, unicode) are handled gracefully."""	159	"""Verify event names with special characters (spaces, symbols, unicode) are handled gracefully."""
160	special_names = ["test event", "test@event", "测试事件", "a" * 256]	160	special_names = ["test event", "test@event", "测试事件", "a" * 256]
161	for name in special_names:	161	for name in special_names:
162	with self.subTest(name=name):	162	with self.subTest(name=name):
163	event = Event(name=name, source=EventSource.WORKER)	163	event = Event(name=name, source=EventSource.WORKER)
164	record(event, destination=self.DESTINATION_NULL)	164	record(event, destination=self.DESTINATION_NULL)
165		165
166	def test_record_all_event_sources(self):	166	def test_record_all_event_sources(self):
167	"""Verify record compatibility with all EventSource enum values."""	167	"""Verify record compatibility with all EventSource enum values."""
168	for source in EventSource:	168	for source in EventSource:
169	with self.subTest(source=source):	169	with self.subTest(source=source):
170	event = Event(	170	event = Event(
171	name=f"{self.TEST_EVENT_PREFIX}source_{source.name.lower()}",	171	name=f"{self.TEST_EVENT_PREFIX}source_{source.name.lower()}",
172	source=source,	172	source=source,
173	metadata={"source_type": source.name}	173	metadata={"source_type": source.name}
174	)	174	)
175	record(event)	175	record(event)
176	self.assertEqual(event.source, source)	176	self.assertEqual(event.source, source)
177	self.assertIn("source_type", event.metadata)	177	self.assertIn("source_type", event.metadata)
178	self.assertEqual(event.metadata["source_type"], source.name)	178	self.assertEqual(event.metadata["source_type"], source.name)
179		179
180		180
181	class TestEventMetadataValue(TestCase):	181	class TestEventMetadataValue(TestCase):
182	"""Test torch.distributed.elastic.events.api.EventMetadataValue type alias."""	182	"""Test torch.distributed.elastic.events.api.EventMetadataValue type alias."""
183		183
184	def test_event_metadata_value_is_defined(self):	184	def test_event_metadata_value_is_defined(self):
185	"""Verify that EventMetadataValue is exported and defined."""	185	"""Verify that EventMetadataValue is exported and defined."""
186	self.assertIsNotNone(EventMetadataValue)	186	self.assertIsNotNone(EventMetadataValue)
187	self.assertIs(get_origin(EventMetadataValue), Union)	187	self.assertIs(get_origin(EventMetadataValue), Union)
188		188
189	def test_event_metadata_value_type_structure(self):	189	def test_event_metadata_value_type_structure(self):
190	"""	190	"""
191	Verify that EventMetadataValue is a Union of str, int, float, bool, None.	191	Verify that EventMetadataValue is a Union of str, int, float, bool, None.
192	Expected: Optional[Union[str, int, float, bool]]	192	Expected: Optional[Union[str, int, float, bool]]
193	"""	193	"""
194	origin = get_origin(EventMetadataValue)	194	origin = get_origin(EventMetadataValue)
195	args = get_args(EventMetadataValue)	195	args = get_args(EventMetadataValue)
196		196
197	self.assertIs(origin, Union)	197	self.assertIs(origin, Union)
198	self.assertIn(str, args)	198	self.assertIn(str, args)
199	self.assertIn(int, args)	199	self.assertIn(int, args)
200	self.assertIn(float, args)	200	self.assertIn(float, args)
201	self.assertIn(bool, args)	201	self.assertIn(bool, args)
202	self.assertIn(type(None), args)	202	self.assertIn(type(None), args)
203	self.assertEqual(len(args), 5)	203	self.assertEqual(len(args), 5)
204		204
205	def test_event_metadata_value_legal_primitives(self):	205	def test_event_metadata_value_legal_primitives(self):
206	"""Verify that all legal primitive values are accepted by Event metadata."""	206	"""Verify that all legal primitive values are accepted by Event metadata."""
207	test_cases = [	207	test_cases = [
208	("string_val", "hello"),	208	("string_val", "hello"),
209	("int_val", 42),	209	("int_val", 42),
210	("int_val_neg", -7),	210	("int_val_neg", -7),
211	("float_val", 3.14159),	211	("float_val", 3.14159),
212	("float_val_zero", 0.0),	212	("float_val_zero", 0.0),
213	("bool_true", True),	213	("bool_true", True),
214	("bool_false", False),	214	("bool_false", False),
215	("none_val", None),	215	("none_val", None),
216	]	216	]
217		217
218	for key, value in test_cases:	218	for key, value in test_cases:
219	with self.subTest(key=key, value=value):	219	with self.subTest(key=key, value=value):
220	event = Event(	220	event = Event(
221	name=f"test_metadata_{key}",	221	name=f"test_metadata_{key}",
222	source=EventSource.WORKER,	222	source=EventSource.WORKER,
223	metadata={key: value}	223	metadata={key: value}
224	)	224	)
225	self.assertIn(key, event.metadata)	225	self.assertIn(key, event.metadata)
226	self.assertEqual(event.metadata[key], value)	226	self.assertEqual(event.metadata[key], value)
227	record(event, destination=TestEventsRecord.DESTINATION_NULL)	227	record(event, destination=TestEventsRecord.DESTINATION_NULL)
228		228
229	def test_event_metadata_value_none_explicit(self):	229	def test_event_metadata_value_none_explicit(self):
230	"""Verify that None is explicitly allowed as a metadata value."""	230	"""Verify that None is explicitly allowed as a metadata value."""
231	event = Event(	231	event = Event(
232	name="test_metadata_none",	232	name="test_metadata_none",
233	source=EventSource.AGENT,	233	source=EventSource.AGENT,
234	metadata={"explicit_none": None}	234	metadata={"explicit_none": None}
235	)	235	)
236	self.assertIsNone(event.metadata["explicit_none"])	236	self.assertIsNone(event.metadata["explicit_none"])
237	record(event, destination=TestEventsRecord.DESTINATION_NULL)	237	record(event, destination=TestEventsRecord.DESTINATION_NULL)
238		238
239	def test_event_metadata_value_mixed_dict(self):	239	def test_event_metadata_value_mixed_dict(self):
240	"""Verify that a metadata dict containing all legal types can be constructed and recorded."""	240	"""Verify that a metadata dict containing all legal types can be constructed and recorded."""
241	metadata = {	241	metadata = {
242	"epoch": 10,	242	"epoch": 10,
243	"loss": 0.1234,	243	"loss": 0.1234,
244	"model_name": "resnet50",	244	"model_name": "resnet50",
245	"is_training": True,	245	"is_training": True,
246	"checkpoint_path": None,	246	"checkpoint_path": None,
247	}	247	}
248	event = Event(	248	event = Event(
249	name="test_metadata_mixed",	249	name="test_metadata_mixed",
250	source=EventSource.WORKER,	250	source=EventSource.WORKER,
251	metadata=metadata	251	metadata=metadata
252	)	252	)
253	self.assertEqual(len(event.metadata), 5)	253	self.assertEqual(len(event.metadata), 5)
254	record(event, destination=TestEventsRecord.DESTINATION_NULL)	254	record(event, destination=TestEventsRecord.DESTINATION_NULL)
255		255
256	def test_event_metadata_value_serialization_roundtrip(self):	256	def test_event_metadata_value_serialization_roundtrip(self):
257	"""	257	"""
258	Verify that metadata values survive the Event.serialize() roundtrip	258	Verify that metadata values survive the Event.serialize() roundtrip
259	and remain as their original Python types.	259	and remain as their original Python types.
260	"""	260	"""
261	metadata = {	261	metadata = {
262	"lr": 0.01,	262	"lr": 0.01,
263	"step": 100,	263	"step": 100,
264	"tag": "train",	264	"tag": "train",
265	"enabled": False,	265	"enabled": False,
266	"optional": None,	266	"optional": None,
267	}	267	}
268	event = Event(	268	event = Event(
269	name="test_serialization",	269	name="test_serialization",
270	source=EventSource.AGENT,	270	source=EventSource.AGENT,
271	metadata=metadata	271	metadata=metadata
272	)	272	)
273	serialized = event.serialize()	273	serialized = event.serialize()
274		274
275	self.assertIsInstance(serialized, str)	275	self.assertIsInstance(serialized, str)
276	data = json.loads(serialized)	276	data = json.loads(serialized)
277	self.assertIn("metadata", data)	277	self.assertIn("metadata", data)
278		278
279	serialized_metadata = data["metadata"]	279	serialized_metadata = data["metadata"]
280	self.assertEqual(serialized_metadata["lr"], 0.01)	280	self.assertEqual(serialized_metadata["lr"], 0.01)
281	self.assertEqual(serialized_metadata["step"], 100)	281	self.assertEqual(serialized_metadata["step"], 100)
282	self.assertEqual(serialized_metadata["tag"], "train")	282	self.assertEqual(serialized_metadata["tag"], "train")
283	self.assertEqual(serialized_metadata["enabled"], False)	283	self.assertEqual(serialized_metadata["enabled"], False)
284	self.assertIsNone(serialized_metadata["optional"])	284	self.assertIsNone(serialized_metadata["optional"])
285		285
286	def test_event_serialize_empty_metadata(self):	286	def test_event_serialize_empty_metadata(self):
287	"""Verify that empty metadata serializes correctly without data loss."""	287	"""Verify that empty metadata serializes correctly without data loss."""
288	event = Event(	288	event = Event(
289	name="test_empty_metadata",	289	name="test_empty_metadata",
290	source=EventSource.WORKER,	290	source=EventSource.WORKER,
291	metadata={}	291	metadata={}
292	)	292	)
293	data = json.loads(event.serialize())	293	data = json.loads(event.serialize())
294	self.assertEqual(data["metadata"], {})	294	self.assertEqual(data["metadata"], {})
295		295
296	def test_event_serialize_unsupported_metadata_type(self):	296	def test_event_serialize_unsupported_metadata_type(self):
297	"""	297	"""
298	Verify that unsupported metadata types (e.g. set, bytes) raise TypeError	298	Verify that unsupported metadata types (e.g. set, bytes) raise TypeError
299	during serialization / record, since Event itself does not validate at construction time.	299	during serialization / record, since Event itself does not validate at construction time.
300	"""	300	"""
301	invalid_metadata_cases = [	301	invalid_metadata_cases = [
302	("set_val", {1, 2, 3}),	302	("set_val", {1, 2, 3}),
303	("bytes_val", b"test"),	303	("bytes_val", b"test"),
304	("object_val", MagicMock()),	304	("object_val", MagicMock()),
305	]	305	]
306		306
307	for key, value in invalid_metadata_cases:	307	for key, value in invalid_metadata_cases:
308	with self.subTest(key=key):	308	with self.subTest(key=key):
309	event = Event(	309	event = Event(
310	name=f"test_invalid_{key}",	310	name=f"test_invalid_{key}",
311	source=EventSource.WORKER,	311	source=EventSource.WORKER,
312	metadata={key: value}	312	metadata={key: value}
313	)	313	)
314	# Event construction succeeds, but serialize/record should fail	314	# Event construction succeeds, but serialize/record should fail
315	with self.assertRaises(TypeError):	315	with self.assertRaises(TypeError):
316	event.serialize()	316	event.serialize()
317		317
318		318
319	if __name__ == "__main__":	319	if __name__ == "__main__":
320	run_tests()	320	run_tests()

Mtest/distributed/rpc/multiprocessing/test_reduction_new.py+0-1

	@@ -127,4 +127,3 @@ class TestReduction(TestCase):
127		127
128	if __name__ == "__main__":	128	if __name__ == "__main__":
129	run_tests()	129	run_tests()
130

Mtest/distributed/test_all_to_all_single.py+1-0

	@@ -133,6 +133,7 @@ class HcclAlltoAllSingleTest(TestCase):
133	HcclAlltoAllSingleTest._test_alltoall_single_2p,	133	HcclAlltoAllSingleTest._test_alltoall_single_2p,
134	HcclAlltoAllSingleTest._init_dist_hccl)	134	HcclAlltoAllSingleTest._init_dist_hccl)
135		135
		136	@unittest.skip("Disabled during A1 to A2 chip transition")
136	@skipIfUnsupportMultiNPU(2)	137	@skipIfUnsupportMultiNPU(2)
137	def test_alltoall_single_2p_size_dist(self):	138	def test_alltoall_single_2p_size_dist(self):
138	self._test_multiprocess_2p(	139	self._test_multiprocess_2p(

Mtest/distributed/test_device_mesh.py+1100-1100

Mtest/distributed/watchdog/watchdog_quick_exit.py+27-27

	@@ -1,27 +1,27 @@
1	import os	1	import os
2	import time	2	import time
3	import datetime	3	import datetime
4	import torch.distributed as dist	4	import torch.distributed as dist
5	import torch	5	import torch
6	import torch_npu	6	import torch_npu
7		7
8		8
9	def main():	9	def main():
10	torch.npu.set_compile_mode(jit_compile=True)	10	torch.npu.set_compile_mode(jit_compile=True)
11	rank = int(os.environ['RANK'])	11	rank = int(os.environ['RANK'])
12	local_rank = int(os.environ['LOCAL_RANK'])	12	local_rank = int(os.environ['LOCAL_RANK'])
13	device = torch.device('npu:{}'.format(local_rank))	13	device = torch.device('npu:{}'.format(local_rank))
14	torch.npu.set_device(device)	14	torch.npu.set_device(device)
15	dist.init_process_group(backend='hccl', rank=rank, world_size=2, timeout=datetime.timedelta(seconds=120))	15	dist.init_process_group(backend='hccl', rank=rank, world_size=2, timeout=datetime.timedelta(seconds=120))
16	tensor = torch.tensor(1).npu()	16	tensor = torch.tensor(1).npu()
17	dist.all_reduce(tensor)	17	dist.all_reduce(tensor)
18	if rank == 0:	18	if rank == 0:
19	x1 = torch.randn(3).float().npu()	19	x1 = torch.randn(3).float().npu()
20	x2 = torch.randn(1).long().npu()	20	x2 = torch.randn(1).long().npu()
21	x3 = torch.randn(1).float().npu()	21	x3 = torch.randn(1).float().npu()
22	y = torch.addcmul(x1, x2, x3)	22	y = torch.addcmul(x1, x2, x3)
23	dist.all_reduce(tensor)	23	dist.all_reduce(tensor)
24		24
25		25
26	if __name__ == "__main__":	26	if __name__ == "__main__":
27	main()	27	main()

Mtest/dynamo/test_stream.py+32-32

	@@ -1,32 +1,32 @@
1	# Owner(s): ["module: dynamo"]	1	# Owner(s): ["module: dynamo"]
2	import functools	2	import functools
3	import unittest	3	import unittest
4	import torch	4	import torch
5	import torch._dynamo.test_case	5	import torch._dynamo.test_case
6	import torch_npu	6	import torch_npu
7		7
8	requires_npu = functools.partial(unittest.skipIf, not torch.npu.is_available(), "requires npu")	8	requires_npu = functools.partial(unittest.skipIf, not torch.npu.is_available(), "requires npu")
9		9
10		10
11	class StreamintoDynamoTests(torch._dynamo.test_case.TestCase):	11	class StreamintoDynamoTests(torch._dynamo.test_case.TestCase):
12		12
13	@requires_npu()	13	@requires_npu()
14	def test_stream(self):	14	def test_stream(self):
15	def model_1(x):	15	def model_1(x):
16	a = x * x	16	a = x * x
17	s = torch.npu.Stream()	17	s = torch.npu.Stream()
18	s.wait_stream(torch.npu.current_stream())	18	s.wait_stream(torch.npu.current_stream())
19	with torch.npu.stream(s):	19	with torch.npu.stream(s):
20	b = x + a	20	b = x + a
21	return b	21	return b
22	inp = torch.randn(2, 8).npu()	22	inp = torch.randn(2, 8).npu()
23	m = torch.compile(model_1, backend="aot_eager", fullgraph=True)	23	m = torch.compile(model_1, backend="aot_eager", fullgraph=True)
24	output = m(inp)	24	output = m(inp)
25	output1 = model_1(inp)	25	output1 = model_1(inp)
26	torch.allclose(output, output1)	26	torch.allclose(output, output1)
27		27
28		28
29	if __name__ == "__main__":	29	if __name__ == "__main__":
30	from torch._dynamo.test_case import run_tests	30	from torch._dynamo.test_case import run_tests
31		31
32	run_tests()	32	run_tests()

Mtest/get_failed_ut_from_log.py+0-1

	@@ -44,4 +44,3 @@ if __name__ == "__main__":
44	args = parser.parse_args()	44	args = parser.parse_args()
45	failed_ut = get_error_or_fail_ut(args.file)	45	failed_ut = get_error_or_fail_ut(args.file)
46	write_to_json(ut_list=failed_ut)	46	write_to_json(ut_list=failed_ut)
47

Mtest/nn/test_uninitialized_parameter_cls_to_become.py+35-35

	@@ -1,36 +1,36 @@
1	import torch	1	import torch
2	from torch.testing._internal.common_utils import TestCase, run_tests	2	from torch.testing._internal.common_utils import TestCase, run_tests
3	import torch_npu	3	import torch_npu
4		4
5	# 关闭NPU JIT编译，减少CI耗时	5	# 关闭NPU JIT编译，减少CI耗时
6	torch_npu.npu.set_compile_mode(jit_compile=False)	6	torch_npu.npu.set_compile_mode(jit_compile=False)
7		7
8		8
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未初始化参数
22	uninit_param = torch.nn.parameter.UninitializedParameter(device="npu")	22	uninit_param = torch.nn.parameter.UninitializedParameter(device="npu")
23	# 2. 绑定自定义类	23	# 2. 绑定自定义类
24	uninit_param.cls_to_become = CustomParameter	24	uninit_param.cls_to_become = CustomParameter
25	# 3. 实例化参数	25	# 3. 实例化参数
26	uninit_param.materialize(shape=(3, 3))	26	uninit_param.materialize(shape=(3, 3))
27		27
28	# 核心断言（全部通过，无报错）	28	# 核心断言（全部通过，无报错）
29	self.assertEqual(uninit_param.shape, torch.Size((3, 3)))	29	self.assertEqual(uninit_param.shape, torch.Size((3, 3)))
30	self.assertEqual(uninit_param.device.type, "npu")	30	self.assertEqual(uninit_param.device.type, "npu")
31	self.assertIsInstance(uninit_param, CustomParameter)	31	self.assertIsInstance(uninit_param, CustomParameter)
32	self.assertEqual(uninit_param.custom_attr, "custom_param") # 现在能正常访问	32	self.assertEqual(uninit_param.custom_attr, "custom_param") # 现在能正常访问
33		33
34		34
35	if __name__ == "__main__":	35	if __name__ == "__main__":
36	run_tests()	36	run_tests()

Mtest/npu/test_amp.py+526-526

	@@ -1,526 +1,526 @@
1	import unittest	1	import unittest
2	from itertools import chain	2	from itertools import chain
3		3
4	import torch	4	import torch
5		5
6	import torch_npu	6	import torch_npu
7	from torch_npu.npu.amp import GradScaler, autocast	7	from torch_npu.npu.amp import GradScaler, autocast
8	from torch_npu.testing.common_utils import SupportedDevices	8	from torch_npu.testing.common_utils import SupportedDevices
9	from torch_npu.testing.testcase import TestCase, run_tests	9	from torch_npu.testing.testcase import TestCase, run_tests
10		10
11		11
12	def make_device_overflow_1():	12	def make_device_overflow_1():
13	float_tensor = torch.tensor([40000.0], dtype=torch.float16).npu()	13	float_tensor = torch.tensor([40000.0], dtype=torch.float16).npu()
14	float_tensor = float_tensor + float_tensor	14	float_tensor = float_tensor + float_tensor
15		15
16		16
17	def make_device_overflow_2(model):	17	def make_device_overflow_2(model):
18	for param in model.parameters():	18	for param in model.parameters():
19	if param.grad is not None:	19	if param.grad is not None:
20	param.grad = torch.full_like(param.grad, float("inf"))	20	param.grad = torch.full_like(param.grad, float("inf"))
21	break	21	break
22		22
23		23
24	class TestAmp(TestCase):	24	class TestAmp(TestCase):
25		25
26	def test_grad_scaling_scale(self):	26	def test_grad_scaling_scale(self):
27	scaler = GradScaler(init_scale=2.)	27	scaler = GradScaler(init_scale=2.)
28	t0 = torch.full((1,), 4.0, dtype=torch.float32, device="npu")	28	t0 = torch.full((1,), 4.0, dtype=torch.float32, device="npu")
29	t1 = torch.full((1,), 4.0, dtype=torch.float32, device="npu")	29	t1 = torch.full((1,), 4.0, dtype=torch.float32, device="npu")
30	# Create some nested iterables of tensors on different devices.	30	# Create some nested iterables of tensors on different devices.
31	outputs = (t1.clone(), (t0.clone(), t1.clone()), [t0.clone(), (t1.clone(), t0.clone())])	31	outputs = (t1.clone(), (t0.clone(), t1.clone()), [t0.clone(), (t1.clone(), t0.clone())])
32	outputs = scaler.scale(outputs)	32	outputs = scaler.scale(outputs)
33	self.assertTrue(outputs[0] == 8.0 and outputs[1][0] == 8.0 and outputs[1][1] == 8.0 and	33	self.assertTrue(outputs[0] == 8.0 and outputs[1][0] == 8.0 and outputs[1][1] == 8.0 and
34	outputs[2][0] == 8.0 and outputs[2][1][0] == 8.0 and outputs[2][1][1] == 8.0)	34	outputs[2][0] == 8.0 and outputs[2][1][0] == 8.0 and outputs[2][1][1] == 8.0)
35	self.assertTrue(scaler._scale.device == t1.device)	35	self.assertTrue(scaler._scale.device == t1.device)
36		36
37	def test_grad_scaling_state_dict(self):	37	def test_grad_scaling_state_dict(self):
38	for lazy_init_scale in True, False:	38	for lazy_init_scale in True, False:
39	s0 = GradScaler(init_scale=3., growth_factor=4., backoff_factor=.5, growth_interval=2)	39	s0 = GradScaler(init_scale=3., growth_factor=4., backoff_factor=.5, growth_interval=2)
40	s1 = GradScaler(init_scale=6., growth_factor=7., backoff_factor=.8, growth_interval=1)	40	s1 = GradScaler(init_scale=6., growth_factor=7., backoff_factor=.8, growth_interval=1)
41		41
42	# sets a random value for load_state_dict to overwrite	42	# sets a random value for load_state_dict to overwrite
43	s1._init_growth_tracker = 7	43	s1._init_growth_tracker = 7
44		44
45	if lazy_init_scale:	45	if lazy_init_scale:
46	# Dummy scale() call to ensure the scale tensor is lazily initialized.	46	# Dummy scale() call to ensure the scale tensor is lazily initialized.
47	s1.scale(torch.full((1,), 4.0, dtype=torch.float32, device="npu"))	47	s1.scale(torch.full((1,), 4.0, dtype=torch.float32, device="npu"))
48	self.assertTrue(isinstance(s1._scale, torch.npu.FloatTensor))	48	self.assertTrue(isinstance(s1._scale, torch.npu.FloatTensor))
49		49
50	s1.load_state_dict(s0.state_dict())	50	s1.load_state_dict(s0.state_dict())
51		51
52	self.assertTrue(s1.get_scale() == 3.)	52	self.assertTrue(s1.get_scale() == 3.)
53	self.assertTrue(s1.get_growth_factor() == 4.)	53	self.assertTrue(s1.get_growth_factor() == 4.)
54	self.assertTrue(s1.get_backoff_factor() == .5)	54	self.assertTrue(s1.get_backoff_factor() == .5)
55	self.assertTrue(s1.get_growth_interval() == 2)	55	self.assertTrue(s1.get_growth_interval() == 2)
56	self.assertTrue(s1._init_growth_tracker == 0)	56	self.assertTrue(s1._init_growth_tracker == 0)
57		57
58	def _create_scaling_models_optimizers(self, device="npu"):	58	def _create_scaling_models_optimizers(self, device="npu"):
59	# Create a module+optimizer that will use scaling, and a control module+optimizer	59	# Create a module+optimizer that will use scaling, and a control module+optimizer
60	# that will not use scaling, against which the scaling-enabled module+optimizer can be compared.	60	# that will not use scaling, against which the scaling-enabled module+optimizer can be compared.
61	mod_control = torch.nn.Sequential(torch.nn.Linear(8, 8), torch.nn.Linear(8, 8)).to(device=device)	61	mod_control = torch.nn.Sequential(torch.nn.Linear(8, 8), torch.nn.Linear(8, 8)).to(device=device)
62	mod_scaling = torch.nn.Sequential(torch.nn.Linear(8, 8), torch.nn.Linear(8, 8)).to(device=device)	62	mod_scaling = torch.nn.Sequential(torch.nn.Linear(8, 8), torch.nn.Linear(8, 8)).to(device=device)
63	for c, s in zip(mod_control.parameters(), mod_scaling.parameters()):	63	for c, s in zip(mod_control.parameters(), mod_scaling.parameters()):
64	s.data.copy_(c.data)	64	s.data.copy_(c.data)
65		65
66	opt_control = torch.optim.SGD(mod_control.parameters(), lr=1.0)	66	opt_control = torch.optim.SGD(mod_control.parameters(), lr=1.0)
67	opt_scaling = torch.optim.SGD(mod_scaling.parameters(), lr=1.0)	67	opt_scaling = torch.optim.SGD(mod_scaling.parameters(), lr=1.0)
68		68
69	ret = (mod_control, mod_scaling, opt_control, opt_scaling)	69	ret = (mod_control, mod_scaling, opt_control, opt_scaling)
70	return ret	70	return ret
71		71
72	def _create_scaling_case(self, device="npu", dtype=torch.float):	72	def _create_scaling_case(self, device="npu", dtype=torch.float):
73	data = [(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device)),	73	data = [(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device)),
74	(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device)),	74	(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device)),
75	(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device)),	75	(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device)),
76	(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device))]	76	(torch.randn((8, 8), dtype=dtype, device=device), torch.randn((8, 8), dtype=dtype, device=device))]
77		77
78	loss_fn = torch.nn.MSELoss().npu()	78	loss_fn = torch.nn.MSELoss().npu()
79		79
80	skip_iter = 2	80	skip_iter = 2
81		81
82	return self._create_scaling_models_optimizers(device=device) + (data, loss_fn, skip_iter)	82	return self._create_scaling_models_optimizers(device=device) + (data, loss_fn, skip_iter)
83		83
84	# _run_scaling_case generalizes some single-optimizer test logic to avoid too much copy-pasting below.	84	# _run_scaling_case generalizes some single-optimizer test logic to avoid too much copy-pasting below.
85	def _run_scaling_case(self, run, unskipped, skipped, atol=1e-7):	85	def _run_scaling_case(self, run, unskipped, skipped, atol=1e-7):
86	# Ensure scaling can be disabled without changing user control flow.	86	# Ensure scaling can be disabled without changing user control flow.
87	for enabled in True, False:	87	for enabled in True, False:
88	mod_control, mod_scaling, opt_control, opt_scaling, data, loss_fn, skip_iter = self._create_scaling_case()	88	mod_control, mod_scaling, opt_control, opt_scaling, data, loss_fn, skip_iter = self._create_scaling_case()
89		89
90	# For functionality, test with a modest initial scale, and an unrealistically-large growth factor	90	# For functionality, test with a modest initial scale, and an unrealistically-large growth factor
91	# so any potential errors with the growth factor handling will be magnified.	91	# so any potential errors with the growth factor handling will be magnified.
92	scaler = GradScaler(init_scale=128., growth_factor=2.0, enabled=enabled, growth_interval=1)	92	scaler = GradScaler(init_scale=128., growth_factor=2.0, enabled=enabled, growth_interval=1)
93		93
94	_ = run(data, mod_control, opt_control, scaler, loss_fn, skip_iter, False)	94	_ = run(data, mod_control, opt_control, scaler, loss_fn, skip_iter, False)
95	ret = run(data, mod_scaling, opt_scaling, scaler, loss_fn, skip_iter, True)	95	ret = run(data, mod_scaling, opt_scaling, scaler, loss_fn, skip_iter, True)
96		96
97	# Allows run() to optionally return a different scaler instance.	97	# Allows run() to optionally return a different scaler instance.
98	scaler = ret if ret else scaler	98	scaler = ret if ret else scaler
99		99
100	# If scaling was enabled, the scale factor should have been multiplied by the growth factor	100	# If scaling was enabled, the scale factor should have been multiplied by the growth factor
101	# len(data) - skipped times and the backoff factor "skipped" times.	101	# len(data) - skipped times and the backoff factor "skipped" times.
102	if enabled:	102	if enabled:
103	net_growth = scaler.get_growth_factor()**unskipped if unskipped > 0 else 1.0	103	net_growth = scaler.get_growth_factor()**unskipped if unskipped > 0 else 1.0
104	net_backoff = scaler.get_backoff_factor()**skipped if skipped > 0 else 1.0	104	net_backoff = scaler.get_backoff_factor()**skipped if skipped > 0 else 1.0
105	self.assertTrue(scaler.get_scale() == (128. * net_growth * net_backoff))	105	self.assertTrue(scaler.get_scale() == (128. * net_growth * net_backoff))
106	else:	106	else:
107	self.assertTrue(scaler.get_scale() == 1.0)	107	self.assertTrue(scaler.get_scale() == 1.0)
108		108
109	for c, s in zip(mod_control.parameters(), mod_scaling.parameters()):	109	for c, s in zip(mod_control.parameters(), mod_scaling.parameters()):
110	c = c.cpu().to(torch.float).detach().numpy()	110	c = c.cpu().to(torch.float).detach().numpy()
111	s = s.cpu().to(torch.float).detach().numpy()	111	s = s.cpu().to(torch.float).detach().numpy()
112	self.assertRtolEqual(c, s, atol)	112	self.assertRtolEqual(c, s, atol)
113		113
114	# Compares no scaling + no autocasting against scaling + autocasting.	114	# Compares no scaling + no autocasting against scaling + autocasting.
115	@SupportedDevices(['Ascend910A', 'Ascend910P'])	115	@SupportedDevices(['Ascend910A', 'Ascend910P'])
116	def test_grad_scaling_autocast_1(self):	116	def test_grad_scaling_autocast_1(self):
117	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	117	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
118	for i, (input_data, target) in enumerate(data):	118	for i, (input_data, target) in enumerate(data):
119	optimizer.zero_grad()	119	optimizer.zero_grad()
120	with torch.autocast('npu', enabled=try_scaling_api):	120	with torch.autocast('npu', enabled=try_scaling_api):
121	output = model(input_data)	121	output = model(input_data)
122	loss = loss_fn(output, target)	122	loss = loss_fn(output, target)
123	if try_scaling_api:	123	if try_scaling_api:
124	scaler.scale(loss).backward()	124	scaler.scale(loss).backward()
125	if i == skip_iter and scaler.is_enabled():	125	if i == skip_iter and scaler.is_enabled():
126	make_device_overflow_1()	126	make_device_overflow_1()
127	scaler.step(optimizer)	127	scaler.step(optimizer)
128	scaler.update()	128	scaler.update()
129	else:	129	else:
130	loss.backward()	130	loss.backward()
131	if (not scaler.is_enabled()) or (i != skip_iter):	131	if (not scaler.is_enabled()) or (i != skip_iter):
132	optimizer.step()	132	optimizer.step()
133	return scaler	133	return scaler
134		134
135	# sets atol=1e-3 because we're comparing pure fp32 arithmetic vs a mixture of fp16 and fp32	135	# sets atol=1e-3 because we're comparing pure fp32 arithmetic vs a mixture of fp16 and fp32
136	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-3)	136	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-3)
137		137
138	@SupportedDevices(['Ascend910B'])	138	@SupportedDevices(['Ascend910B'])
139	def test_grad_scaling_autocast_2(self):	139	def test_grad_scaling_autocast_2(self):
140	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	140	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
141	for i, (input_data, target) in enumerate(data):	141	for i, (input_data, target) in enumerate(data):
142	optimizer.zero_grad()	142	optimizer.zero_grad()
143	with torch.autocast('npu', enabled=try_scaling_api):	143	with torch.autocast('npu', enabled=try_scaling_api):
144	output = model(input_data)	144	output = model(input_data)
145	loss = loss_fn(output, target)	145	loss = loss_fn(output, target)
146	if try_scaling_api:	146	if try_scaling_api:
147	scaler.scale(loss).backward()	147	scaler.scale(loss).backward()
148	if i == skip_iter and scaler.is_enabled():	148	if i == skip_iter and scaler.is_enabled():
149	make_device_overflow_2(model)	149	make_device_overflow_2(model)
150	scaler.step(optimizer)	150	scaler.step(optimizer)
151	scaler.update()	151	scaler.update()
152	else:	152	else:
153	loss.backward()	153	loss.backward()
154	if (not scaler.is_enabled()) or (i != skip_iter):	154	if (not scaler.is_enabled()) or (i != skip_iter):
155	optimizer.step()	155	optimizer.step()
156	return scaler	156	return scaler
157		157
158	# sets atol=1e-3 because we're comparing pure fp32 arithmetic vs a mixture of fp16 and fp32	158	# sets atol=1e-3 because we're comparing pure fp32 arithmetic vs a mixture of fp16 and fp32
159	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-3)	159	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-3)
160		160
161	@SupportedDevices(['Ascend910A', 'Ascend910P'])	161	@SupportedDevices(['Ascend910A', 'Ascend910P'])
162	def test_grad_scaling_clipping_1(self):	162	def test_grad_scaling_clipping_1(self):
163	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	163	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
164	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads	164	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads
165	for i, (input_data, target) in enumerate(data):	165	for i, (input_data, target) in enumerate(data):
166	optimizer.zero_grad()	166	optimizer.zero_grad()
167	output = model(input_data)	167	output = model(input_data)
168	loss = loss_fn(output, target)	168	loss = loss_fn(output, target)
169	if try_scaling_api:	169	if try_scaling_api:
170	scaler.scale(loss).backward()	170	scaler.scale(loss).backward()
171	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm * scaler.get_scale())	171	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm * scaler.get_scale())
172	if i == skip_iter and scaler.is_enabled():	172	if i == skip_iter and scaler.is_enabled():
173	make_device_overflow_1()	173	make_device_overflow_1()
174	scaler.step(optimizer)	174	scaler.step(optimizer)
175	scaler.update()	175	scaler.update()
176	else:	176	else:
177	loss.backward()	177	loss.backward()
178	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)	178	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)
179	if (not scaler.is_enabled()) or (i != skip_iter):	179	if (not scaler.is_enabled()) or (i != skip_iter):
180	optimizer.step()	180	optimizer.step()
181		181
182	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-6)	182	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-6)
183		183
184	@SupportedDevices(['Ascend910B'])	184	@SupportedDevices(['Ascend910B'])
185	def test_grad_scaling_clipping_2(self):	185	def test_grad_scaling_clipping_2(self):
186	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	186	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
187	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads	187	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads
188	for i, (input_data, target) in enumerate(data):	188	for i, (input_data, target) in enumerate(data):
189	optimizer.zero_grad()	189	optimizer.zero_grad()
190	output = model(input_data)	190	output = model(input_data)
191	loss = loss_fn(output, target)	191	loss = loss_fn(output, target)
192	if try_scaling_api:	192	if try_scaling_api:
193	scaler.scale(loss).backward()	193	scaler.scale(loss).backward()
194	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm * scaler.get_scale())	194	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm * scaler.get_scale())
195	if i == skip_iter and scaler.is_enabled():	195	if i == skip_iter and scaler.is_enabled():
196	make_device_overflow_2(model)	196	make_device_overflow_2(model)
197	scaler.step(optimizer)	197	scaler.step(optimizer)
198	scaler.update()	198	scaler.update()
199	else:	199	else:
200	loss.backward()	200	loss.backward()
201	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)	201	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)
202	if (not scaler.is_enabled()) or (i != skip_iter):	202	if (not scaler.is_enabled()) or (i != skip_iter):
203	optimizer.step()	203	optimizer.step()
204		204
205	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-6)	205	self._run_scaling_case(run, unskipped=3, skipped=1, atol=1e-6)
206		206
207	@SupportedDevices(['Ascend910A', 'Ascend910P'])	207	@SupportedDevices(['Ascend910A', 'Ascend910P'])
208	def test_grad_scaling_clipping_separate_unscale_1(self):	208	def test_grad_scaling_clipping_separate_unscale_1(self):
209	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	209	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
210	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads	210	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads
211	for i, (input_data, target) in enumerate(data):	211	for i, (input_data, target) in enumerate(data):
212	optimizer.zero_grad()	212	optimizer.zero_grad()
213	output = model(input_data)	213	output = model(input_data)
214	loss = loss_fn(output, target)	214	loss = loss_fn(output, target)
215	if try_scaling_api:	215	if try_scaling_api:
216	scaler.scale(loss).backward()	216	scaler.scale(loss).backward()
217	if i == skip_iter and scaler.is_enabled():	217	if i == skip_iter and scaler.is_enabled():
218	make_device_overflow_1()	218	make_device_overflow_1()
219	scaler.unscale_(optimizer)	219	scaler.unscale_(optimizer)
220	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)	220	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)
221	scaler.step(optimizer)	221	scaler.step(optimizer)
222	scaler.update()	222	scaler.update()
223	else:	223	else:
224	loss.backward()	224	loss.backward()
225	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)	225	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)
226	if (not scaler.is_enabled()) or (i != skip_iter):	226	if (not scaler.is_enabled()) or (i != skip_iter):
227	optimizer.step()	227	optimizer.step()
228		228
229	self._run_scaling_case(run, unskipped=3, skipped=1)	229	self._run_scaling_case(run, unskipped=3, skipped=1)
230		230
231	@SupportedDevices(['Ascend910B'])	231	@SupportedDevices(['Ascend910B'])
232	def test_grad_scaling_clipping_separate_unscale_2(self):	232	def test_grad_scaling_clipping_separate_unscale_2(self):
233	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	233	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
234	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads	234	max_norm = 0.2 # A reasonable value that actually has an effect, based on printouts of grads
235	for i, (input_data, target) in enumerate(data):	235	for i, (input_data, target) in enumerate(data):
236	optimizer.zero_grad()	236	optimizer.zero_grad()
237	output = model(input_data)	237	output = model(input_data)
238	loss = loss_fn(output, target)	238	loss = loss_fn(output, target)
239	if try_scaling_api:	239	if try_scaling_api:
240	scaler.scale(loss).backward()	240	scaler.scale(loss).backward()
241	if i == skip_iter and scaler.is_enabled():	241	if i == skip_iter and scaler.is_enabled():
242	make_device_overflow_2(model)	242	make_device_overflow_2(model)
243	scaler.unscale_(optimizer)	243	scaler.unscale_(optimizer)
244	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)	244	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)
245	scaler.step(optimizer)	245	scaler.step(optimizer)
246	scaler.update()	246	scaler.update()
247	else:	247	else:
248	loss.backward()	248	loss.backward()
249	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)	249	torch.nn.utils.clip_grad_norm_(model.parameters(), max_norm)
250	if (not scaler.is_enabled()) or (i != skip_iter):	250	if (not scaler.is_enabled()) or (i != skip_iter):
251	optimizer.step()	251	optimizer.step()
252		252
253	self._run_scaling_case(run, unskipped=3, skipped=1)	253	self._run_scaling_case(run, unskipped=3, skipped=1)
254		254
255	@SupportedDevices(['Ascend910A', 'Ascend910P'])	255	@SupportedDevices(['Ascend910A', 'Ascend910P'])
256	def test_grad_scaling_penalty_1(self):	256	def test_grad_scaling_penalty_1(self):
257	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	257	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
258	for i, (input_data, target) in enumerate(data):	258	for i, (input_data, target) in enumerate(data):
259	optimizer.zero_grad()	259	optimizer.zero_grad()
260	output = model(input_data)	260	output = model(input_data)
261	loss = loss_fn(output, target)	261	loss = loss_fn(output, target)
262		262
263	if try_scaling_api:	263	if try_scaling_api:
264	grad_params = torch.autograd.grad(scaler.scale(loss),	264	grad_params = torch.autograd.grad(scaler.scale(loss),
265	model.parameters(), create_graph=True)	265	model.parameters(), create_graph=True)
266	inv_scale = 1. / scaler.get_scale()	266	inv_scale = 1. / scaler.get_scale()
267	grad_params = [p * inv_scale for p in grad_params]	267	grad_params = [p * inv_scale for p in grad_params]
268	else:	268	else:
269	grad_params = torch.autograd.grad(loss, model.parameters(), create_graph=True)	269	grad_params = torch.autograd.grad(loss, model.parameters(), create_graph=True)
270		270
271	grad_norm = 0	271	grad_norm = 0
272	for grad in grad_params:	272	for grad in grad_params:
273	grad_norm += grad.pow(2).sum()	273	grad_norm += grad.pow(2).sum()
274	grad_norm = grad_norm.sqrt()	274	grad_norm = grad_norm.sqrt()
275	loss = loss + grad_norm	275	loss = loss + grad_norm
276		276
277	if try_scaling_api:	277	if try_scaling_api:
278	scaler.scale(loss).backward()	278	scaler.scale(loss).backward()
279	if i == skip_iter and scaler.is_enabled():	279	if i == skip_iter and scaler.is_enabled():
280	make_device_overflow_1()	280	make_device_overflow_1()
281	scaler.step(optimizer)	281	scaler.step(optimizer)
282	scaler.update()	282	scaler.update()
283	else:	283	else:
284	loss.backward()	284	loss.backward()
285	if (not scaler.is_enabled()) or (i != skip_iter):	285	if (not scaler.is_enabled()) or (i != skip_iter):
286	optimizer.step()	286	optimizer.step()
287		287
288	self._run_scaling_case(run, unskipped=3, skipped=1)	288	self._run_scaling_case(run, unskipped=3, skipped=1)
289		289
290	@SupportedDevices(['Ascend910B'])	290	@SupportedDevices(['Ascend910B'])
291	def test_grad_scaling_penalty_2(self):	291	def test_grad_scaling_penalty_2(self):
292	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	292	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
293	for i, (input_data, target) in enumerate(data):	293	for i, (input_data, target) in enumerate(data):
294	optimizer.zero_grad()	294	optimizer.zero_grad()
295	output = model(input_data)	295	output = model(input_data)
296	loss = loss_fn(output, target)	296	loss = loss_fn(output, target)
297		297
298	if try_scaling_api:	298	if try_scaling_api:
299	grad_params = torch.autograd.grad(scaler.scale(loss),	299	grad_params = torch.autograd.grad(scaler.scale(loss),
300	model.parameters(), create_graph=True)	300	model.parameters(), create_graph=True)
301	inv_scale = 1. / scaler.get_scale()	301	inv_scale = 1. / scaler.get_scale()
302	grad_params = [p * inv_scale for p in grad_params]	302	grad_params = [p * inv_scale for p in grad_params]
303	else:	303	else:
304	grad_params = torch.autograd.grad(loss, model.parameters(), create_graph=True)	304	grad_params = torch.autograd.grad(loss, model.parameters(), create_graph=True)
305		305
306	grad_norm = 0	306	grad_norm = 0
307	for grad in grad_params:	307	for grad in grad_params:
308	grad_norm += grad.pow(2).sum()	308	grad_norm += grad.pow(2).sum()
309	grad_norm = grad_norm.sqrt()	309	grad_norm = grad_norm.sqrt()
310	loss = loss + grad_norm	310	loss = loss + grad_norm
311		311
312	if try_scaling_api:	312	if try_scaling_api:
313	scaler.scale(loss).backward()	313	scaler.scale(loss).backward()
314	if i == skip_iter and scaler.is_enabled():	314	if i == skip_iter and scaler.is_enabled():
315	make_device_overflow_2(model)	315	make_device_overflow_2(model)
316	scaler.step(optimizer)	316	scaler.step(optimizer)
317	scaler.update()	317	scaler.update()
318	else:	318	else:
319	loss.backward()	319	loss.backward()
320	if (not scaler.is_enabled()) or (i != skip_iter):	320	if (not scaler.is_enabled()) or (i != skip_iter):
321	optimizer.step()	321	optimizer.step()
322		322
323	self._run_scaling_case(run, unskipped=3, skipped=1)	323	self._run_scaling_case(run, unskipped=3, skipped=1)
324		324
325	def test_grad_scaling_accumulation(self):	325	def test_grad_scaling_accumulation(self):
326	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):	326	def run(data, model, optimizer, scaler, loss_fn, skip_iter, try_scaling_api):
327	iters_to_accumulate = 2	327	iters_to_accumulate = 2
328	for i, (input_data, target) in enumerate(data):	328	for i, (input_data, target) in enumerate(data):
329	output = model(input_data)	329	output = model(input_data)
330	loss = loss_fn(output, target)	330	loss = loss_fn(output, target)
331	loss = loss / iters_to_accumulate	331	loss = loss / iters_to_accumulate
332	if try_scaling_api:	332	if try_scaling_api:
333	scaler.scale(loss).backward()	333	scaler.scale(loss).backward()
334	else:	334	else:
335	loss.backward()	335	loss.backward()
336	if (i + 1) % iters_to_accumulate == 0:	336	if (i + 1) % iters_to_accumulate == 0:
337	if try_scaling_api:	337	if try_scaling_api:
338	scaler.step(optimizer)	338	scaler.step(optimizer)
339	scaler.update()	339	scaler.update()
340	optimizer.zero_grad()	340	optimizer.zero_grad()
341	else:	341	else:
342	optimizer.step()	342	optimizer.step()
343	optimizer.zero_grad()	343	optimizer.zero_grad()
344		344
345	self._run_scaling_case(run, unskipped=2, skipped=0)	345	self._run_scaling_case(run, unskipped=2, skipped=0)
346		346
347	@SupportedDevices(['Ascend910A', 'Ascend910P'])	347	@SupportedDevices(['Ascend910A', 'Ascend910P'])
348	def test_grad_scaling_multiple_1(self):	348	def test_grad_scaling_multiple_1(self):
349	# Tests gradient scaling with 2 models and 2 optimizers that both receive gradients from 2 losses.	349	# Tests gradient scaling with 2 models and 2 optimizers that both receive gradients from 2 losses.
350	# Some of the logic here cannot reuse the generic helper functions created for the 1-optimizer cases.	350	# Some of the logic here cannot reuse the generic helper functions created for the 1-optimizer cases.
351	for enabled in True, False:	351	for enabled in True, False:
352	mod_control0, mod_scaling0, opt_control0, opt_scaling0, data, loss_fn, skip_iter = \	352	mod_control0, mod_scaling0, opt_control0, opt_scaling0, data, loss_fn, skip_iter = \
353	self._create_scaling_case()	353	self._create_scaling_case()
354	mod_control1, mod_scaling1, opt_control1, opt_scaling1 = \	354	mod_control1, mod_scaling1, opt_control1, opt_scaling1 = \
355	self._create_scaling_models_optimizers()	355	self._create_scaling_models_optimizers()
356		356
357	scaler = GradScaler(init_scale=128., growth_factor=2.0, enabled=enabled, growth_interval=1)	357	scaler = GradScaler(init_scale=128., growth_factor=2.0, enabled=enabled, growth_interval=1)
358		358
359	def run(model0, model1, optimizer0, optimizer1, try_scaling_api):	359	def run(model0, model1, optimizer0, optimizer1, try_scaling_api):
360	for i, (input_data, target) in enumerate(data):	360	for i, (input_data, target) in enumerate(data):
361	optimizer0.zero_grad()	361	optimizer0.zero_grad()
362	optimizer1.zero_grad()	362	optimizer1.zero_grad()
363	output0 = model0(input_data)	363	output0 = model0(input_data)
364	output1 = model1(input_data)	364	output1 = model1(input_data)
365	loss0 = loss_fn(0.3 * output0 + 0.7 * output1, target)	365	loss0 = loss_fn(0.3 * output0 + 0.7 * output1, target)
366	loss1 = loss_fn(0.6 * output0 - 0.4 * output1, target)	366	loss1 = loss_fn(0.6 * output0 - 0.4 * output1, target)
367		367
368	if try_scaling_api:	368	if try_scaling_api:
369	scaler.scale(loss0).backward(retain_graph=True)	369	scaler.scale(loss0).backward(retain_graph=True)
370	scaler.scale(loss1).backward()	370	scaler.scale(loss1).backward()
371	if i == skip_iter and scaler.is_enabled():	371	if i == skip_iter and scaler.is_enabled():
372	make_device_overflow_1()	372	make_device_overflow_1()
373		373
374	# As an additional stress test, separately unscale for one of the optimizers.	374	# As an additional stress test, separately unscale for one of the optimizers.
375	scaler.unscale_(optimizer0)	375	scaler.unscale_(optimizer0)
376		376
377	scaler.step(optimizer0)	377	scaler.step(optimizer0)
378	scaler.step(optimizer1)	378	scaler.step(optimizer1)
379	scaler.update()	379	scaler.update()
380	else:	380	else:
381	loss0.backward(retain_graph=True)	381	loss0.backward(retain_graph=True)
382	loss1.backward()	382	loss1.backward()
383	if (not scaler.is_enabled()) or (i != skip_iter):	383	if (not scaler.is_enabled()) or (i != skip_iter):
384	optimizer0.step()	384	optimizer0.step()
385	optimizer1.step()	385	optimizer1.step()
386		386
387	run(mod_control0, mod_control1, opt_control0, opt_control1, False)	387	run(mod_control0, mod_control1, opt_control0, opt_control1, False)
388	run(mod_scaling0, mod_scaling1, opt_scaling0, opt_scaling1, True)	388	run(mod_scaling0, mod_scaling1, opt_scaling0, opt_scaling1, True)
389		389
390	# The loss scale should have been multiplied by the growth factor 3 times and the backoff factor once.	390	# The loss scale should have been multiplied by the growth factor 3 times and the backoff factor once.
391	self.assertTrue(scaler.get_scale() == (128. * scaler.get_growth_factor()*3	391	self.assertTrue(scaler.get_scale() == (128. * scaler.get_growth_factor()*3
392	scaler.get_backoff_factor()**1) if enabled else 1.0)	392	scaler.get_backoff_factor()**1) if enabled else 1.0)
393		393
394	for c, s in zip(chain(mod_control0.parameters(), mod_control1.parameters()),	394	for c, s in zip(chain(mod_control0.parameters(), mod_control1.parameters()),
395	chain(mod_scaling0.parameters(), mod_scaling1.parameters())):	395	chain(mod_scaling0.parameters(), mod_scaling1.parameters())):
396	c = c.cpu().to(torch.float).detach().numpy()	396	c = c.cpu().to(torch.float).detach().numpy()
397	s = s.cpu().to(torch.float).detach().numpy()	397	s = s.cpu().to(torch.float).detach().numpy()
398	self.assertRtolEqual(c, s, 1e-7)	398	self.assertRtolEqual(c, s, 1e-7)
399		399
400	@SupportedDevices(['Ascend910B'])	400	@SupportedDevices(['Ascend910B'])
401	def test_grad_scaling_multiple_2(self):	401	def test_grad_scaling_multiple_2(self):
402	# Tests gradient scaling with 2 models and 2 optimizers that both receive gradients from 2 losses.	402	# Tests gradient scaling with 2 models and 2 optimizers that both receive gradients from 2 losses.
403	# Some of the logic here cannot reuse the generic helper functions created for the 1-optimizer cases.	403	# Some of the logic here cannot reuse the generic helper functions created for the 1-optimizer cases.
404	for enabled in True, False:	404	for enabled in True, False:
405	mod_control0, mod_scaling0, opt_control0, opt_scaling0, data, loss_fn, skip_iter = \	405	mod_control0, mod_scaling0, opt_control0, opt_scaling0, data, loss_fn, skip_iter = \
406	self._create_scaling_case()	406	self._create_scaling_case()
407	mod_control1, mod_scaling1, opt_control1, opt_scaling1 = \	407	mod_control1, mod_scaling1, opt_control1, opt_scaling1 = \
408	self._create_scaling_models_optimizers()	408	self._create_scaling_models_optimizers()
409		409
410	scaler = GradScaler(init_scale=128., growth_factor=2.0, enabled=enabled, growth_interval=1)	410	scaler = GradScaler(init_scale=128., growth_factor=2.0, enabled=enabled, growth_interval=1)
411		411
412	def run(model0, model1, optimizer0, optimizer1, try_scaling_api):	412	def run(model0, model1, optimizer0, optimizer1, try_scaling_api):
413	for i, (input_data, target) in enumerate(data):	413	for i, (input_data, target) in enumerate(data):
414	optimizer0.zero_grad()	414	optimizer0.zero_grad()
415	optimizer1.zero_grad()	415	optimizer1.zero_grad()
416	output0 = model0(input_data)	416	output0 = model0(input_data)
417	output1 = model1(input_data)	417	output1 = model1(input_data)
418	loss0 = loss_fn(0.3 * output0 + 0.7 * output1, target)	418	loss0 = loss_fn(0.3 * output0 + 0.7 * output1, target)
419	loss1 = loss_fn(0.6 * output0 - 0.4 * output1, target)	419	loss1 = loss_fn(0.6 * output0 - 0.4 * output1, target)
420		420
421	if try_scaling_api:	421	if try_scaling_api:
422	scaler.scale(loss0).backward(retain_graph=True)	422	scaler.scale(loss0).backward(retain_graph=True)
423	scaler.scale(loss1).backward()	423	scaler.scale(loss1).backward()
424	if i == skip_iter and scaler.is_enabled():	424	if i == skip_iter and scaler.is_enabled():
425	make_device_overflow_2(model0)	425	make_device_overflow_2(model0)
426		426
427	# As an additional stress test, separately unscale for one of the optimizers.	427	# As an additional stress test, separately unscale for one of the optimizers.
428	scaler.unscale_(optimizer0)	428	scaler.unscale_(optimizer0)
429		429
430	scaler.step(optimizer0)	430	scaler.step(optimizer0)
431	scaler.step(optimizer1)	431	scaler.step(optimizer1)
432	scaler.update()	432	scaler.update()
433	else:	433	else:
434	loss0.backward(retain_graph=True)	434	loss0.backward(retain_graph=True)
435	loss1.backward()	435	loss1.backward()
436	if (not scaler.is_enabled()) or (i != skip_iter):	436	if (not scaler.is_enabled()) or (i != skip_iter):
437	optimizer0.step()	437	optimizer0.step()
438	optimizer1.step()	438	optimizer1.step()
439		439
440	run(mod_control0, mod_control1, opt_control0, opt_control1, False)	440	run(mod_control0, mod_control1, opt_control0, opt_control1, False)
441	run(mod_scaling0, mod_scaling1, opt_scaling0, opt_scaling1, True)	441	run(mod_scaling0, mod_scaling1, opt_scaling0, opt_scaling1, True)
442		442
443	# The loss scale should have been multiplied by the growth factor 3 times and the backoff factor once.	443	# The loss scale should have been multiplied by the growth factor 3 times and the backoff factor once.
444	self.assertTrue(scaler.get_scale() == (128. * scaler.get_growth_factor() ** 3 *	444	self.assertTrue(scaler.get_scale() == (128. * scaler.get_growth_factor() ** 3 *
445	scaler.get_backoff_factor() ** 1) if enabled else 1.0)	445	scaler.get_backoff_factor() ** 1) if enabled else 1.0)
446		446
447	for c, s in zip(chain(mod_control0.parameters(), mod_control1.parameters()),	447	for c, s in zip(chain(mod_control0.parameters(), mod_control1.parameters()),
448	chain(mod_scaling0.parameters(), mod_scaling1.parameters())):	448	chain(mod_scaling0.parameters(), mod_scaling1.parameters())):
449	c = c.cpu().to(torch.float).detach().numpy()	449	c = c.cpu().to(torch.float).detach().numpy()
450	s = s.cpu().to(torch.float).detach().numpy()	450	s = s.cpu().to(torch.float).detach().numpy()
451	self.assertRtolEqual(c, s, 1e-7)	451	self.assertRtolEqual(c, s, 1e-7)
452		452
453	def test_autocast_custom_enabled(self):	453	def test_autocast_custom_enabled(self):
454	class MyMM(torch.autograd.Function):	454	class MyMM(torch.autograd.Function):
455	@staticmethod	455	@staticmethod
456	@torch.npu.amp.custom_fwd	456	@torch.npu.amp.custom_fwd
457	def forward(ctx, a, b):	457	def forward(ctx, a, b):
458	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))	458	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))
459	self.assertTrue(a.dtype is torch.float32)	459	self.assertTrue(a.dtype is torch.float32)
460	self.assertTrue(b.dtype is torch.float32)	460	self.assertTrue(b.dtype is torch.float32)
461	self.assertTrue(torch.npu.is_autocast_enabled())	461	self.assertTrue(torch.npu.is_autocast_enabled())
462	ctx.save_for_backward(a, b)	462	ctx.save_for_backward(a, b)
463	return a.mm(b)	463	return a.mm(b)
464		464
465	@staticmethod	465	@staticmethod
466	@torch.npu.amp.custom_bwd	466	@torch.npu.amp.custom_bwd
467	def backward(ctx, grad):	467	def backward(ctx, grad):
468	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))	468	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))
469	self.assertTrue(torch.npu.is_autocast_enabled())	469	self.assertTrue(torch.npu.is_autocast_enabled())
470	a, b = ctx.saved_tensors	470	a, b = ctx.saved_tensors
471	return grad.mm(b.t()), a.t().mm(grad)	471	return grad.mm(b.t()), a.t().mm(grad)
472		472
473	mymm = MyMM.apply	473	mymm = MyMM.apply
474		474
475	x = torch.randn((8, 8), device="npu", dtype=torch.float32, requires_grad=True)	475	x = torch.randn((8, 8), device="npu", dtype=torch.float32, requires_grad=True)
476	y = torch.randn((8, 8), device="npu", dtype=torch.float32, requires_grad=True)	476	y = torch.randn((8, 8), device="npu", dtype=torch.float32, requires_grad=True)
477		477
478	with torch.npu.amp.autocast():	478	with torch.npu.amp.autocast():
479	output = mymm(x, y)	479	output = mymm(x, y)
480	self.assertTrue(output.dtype is torch.float16)	480	self.assertTrue(output.dtype is torch.float16)
481	loss = output.sum()	481	loss = output.sum()
482	loss.backward()	482	loss.backward()
483		483
484	def test_autocast_custom_cast_inputs(self):	484	def test_autocast_custom_cast_inputs(self):
485	class MyMM(torch.autograd.Function):	485	class MyMM(torch.autograd.Function):
486	@staticmethod	486	@staticmethod
487	@torch.npu.amp.custom_fwd(cast_inputs=torch.float32)	487	@torch.npu.amp.custom_fwd(cast_inputs=torch.float32)
488	def forward(ctx, a, container, expect_type):	488	def forward(ctx, a, container, expect_type):
489	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))	489	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))
490	b = container[1][0]	490	b = container[1][0]
491	self.assertTrue(a.dtype is expect_type)	491	self.assertTrue(a.dtype is expect_type)
492	self.assertTrue(b.dtype is expect_type)	492	self.assertTrue(b.dtype is expect_type)
493	self.assertFalse(torch.npu.is_autocast_enabled())	493	self.assertFalse(torch.npu.is_autocast_enabled())
494	ctx.save_for_backward(a, b)	494	ctx.save_for_backward(a, b)
495	return a.mm(b)	495	return a.mm(b)
496		496
497	@staticmethod	497	@staticmethod
498	@torch.npu.amp.custom_bwd	498	@torch.npu.amp.custom_bwd
499	def backward(ctx, grad):	499	def backward(ctx, grad):
500	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))	500	self.assertTrue(ctx._dtype is torch.get_autocast_dtype("npu"))
501	a, b = ctx.saved_tensors	501	a, b = ctx.saved_tensors
502	return grad.mm(b.t()), None, None	502	return grad.mm(b.t()), None, None
503		503
504	mymm = MyMM.apply	504	mymm = MyMM.apply
505		505
506	x = torch.randn((8, 8), device="npu", dtype=torch.float16, requires_grad=True)	506	x = torch.randn((8, 8), device="npu", dtype=torch.float16, requires_grad=True)
507	# Puts one input tensor in a nested container. y's contained Tensor won't receive a gradient,	507	# Puts one input tensor in a nested container. y's contained Tensor won't receive a gradient,
508	# because torch.autograd.Function can't hand gradients back to non-Tensor forward arguments.	508	# because torch.autograd.Function can't hand gradients back to non-Tensor forward arguments.
509	# Sets requires_grad=False explicitly so we don't lie about expecting a gradient.	509	# Sets requires_grad=False explicitly so we don't lie about expecting a gradient.
510	y = (0, {0: torch.randn((8, 8), device="npu", dtype=torch.float16, requires_grad=False)})	510	y = (0, {0: torch.randn((8, 8), device="npu", dtype=torch.float16, requires_grad=False)})
511		511
512	with torch.autocast('npu', ):	512	with torch.autocast('npu', ):
513	output = mymm(x, y, torch.float32)	513	output = mymm(x, y, torch.float32)
514	self.assertTrue(output.dtype is torch.float32)	514	self.assertTrue(output.dtype is torch.float32)
515	loss = output.sum()	515	loss = output.sum()
516	loss.backward()	516	loss.backward()
517		517
518	# Tests if custom_fwd becomes a no-op when mymm runs outside an autocast-enabled region.	518	# Tests if custom_fwd becomes a no-op when mymm runs outside an autocast-enabled region.
519	output = mymm(x, y, torch.float16)	519	output = mymm(x, y, torch.float16)
520	self.assertTrue(output.dtype is torch.float16)	520	self.assertTrue(output.dtype is torch.float16)
521	loss = output.sum()	521	loss = output.sum()
522	loss.backward()	522	loss.backward()
523		523
524		524
525	if __name__ == "__main__":	525	if __name__ == "__main__":
526	run_tests()	526	run_tests()

Mtest/npu/test_copy.py+137-137

Mtest/npu/test_expandable_segments.py+105-105

	@@ -1,105 +1,105 @@
1	import os	1	import os
2	import gc	2	import gc
3	import unittest	3	import unittest
4		4
5	import torch	5	import torch
6	import torch_npu	6	import torch_npu
7	from torch_npu.testing.testcase import TestCase, run_tests	7	from torch_npu.testing.testcase import TestCase, run_tests
8	from torch.testing._internal.common_utils import TestCase, run_tests, TEST_PRIVATEUSE1	8	from torch.testing._internal.common_utils import TestCase, run_tests, TEST_PRIVATEUSE1
9		9
10	os.environ["PYTORCH_NPU_ALLOC_CONF"] = "expandable_segments:True"	10	os.environ["PYTORCH_NPU_ALLOC_CONF"] = "expandable_segments:True"
11		11
12	device_name = torch_npu.npu.get_device_name(0)	12	device_name = torch_npu.npu.get_device_name(0)
13		13
14		14
15	class Test_expandable_segments(TestCase):	15	class Test_expandable_segments(TestCase):
16	@unittest.skipIf(device_name == "Ascend910B4", "Skip when device name is Ascend910B4")	16	@unittest.skipIf(device_name == "Ascend910B4", "Skip when device name is Ascend910B4")
17	def test_empty_virt_addr_cache(self):	17	def test_empty_virt_addr_cache(self):
18	gc.collect()	18	gc.collect()
19	torch_npu.npu.empty_cache()	19	torch_npu.npu.empty_cache()
20	prev = 0	20	prev = 0
21		21
22	x = torch.empty((7500, 1024, 1024), device="npu")	22	x = torch.empty((7500, 1024, 1024), device="npu")
23	del x	23	del x
24	last_r = torch_npu.npu.memory_reserved()	24	last_r = torch_npu.npu.memory_reserved()
25		25
26	torch_npu.npu.empty_virt_addr_cache()	26	torch_npu.npu.empty_virt_addr_cache()
27	new_r = torch_npu.npu.memory_reserved()	27	new_r = torch_npu.npu.memory_reserved()
28	self.assertEqual(new_r, prev)	28	self.assertEqual(new_r, prev)
29	self.assertEqual(torch_npu.npu.max_memory_reserved(), last_r)	29	self.assertEqual(torch_npu.npu.max_memory_reserved(), last_r)
30		30
31	# test re-alloc after empty virtual address	31	# test re-alloc after empty virtual address
32	try:	32	try:
33	y = torch.empty((7500, 1024, 1024), device="npu")	33	y = torch.empty((7500, 1024, 1024), device="npu")
34	self.assertGreater(torch_npu.npu.memory_allocated(), prev)	34	self.assertGreater(torch_npu.npu.memory_allocated(), prev)
35	finally:	35	finally:
36	if y is not None:	36	if y is not None:
37	del y	37	del y
38	self.assertEqual(torch_npu.npu.memory_allocated(), prev)	38	self.assertEqual(torch_npu.npu.memory_allocated(), prev)
39	torch_npu.npu.empty_virt_addr_cache()	39	torch_npu.npu.empty_virt_addr_cache()
40	# empty unmapped physical handles with empty_cache()	40	# empty unmapped physical handles with empty_cache()
41	torch_npu.npu.empty_cache()	41	torch_npu.npu.empty_cache()
42	self.assertEqual(torch_npu.npu.memory_reserved(), prev)	42	self.assertEqual(torch_npu.npu.memory_reserved(), prev)
43		43
44	@unittest.skipIf(TEST_PRIVATEUSE1, "NPU not available for graph capture")	44	@unittest.skipIf(TEST_PRIVATEUSE1, "NPU not available for graph capture")
45	def test_set_segment_state_to_checkpoint_when_expandable_segments(self):	45	def test_set_segment_state_to_checkpoint_when_expandable_segments(self):
46	def tensor_metadata(x):	46	def tensor_metadata(x):
47	return {	47	return {
48	"nbytes": x.untyped_storage().nbytes(),	48	"nbytes": x.untyped_storage().nbytes(),
49	"data_ptr": x.untyped_storage().data_ptr(),	49	"data_ptr": x.untyped_storage().data_ptr(),
50	"size": x.shape,	50	"size": x.shape,
51	"stride": x.stride(),	51	"stride": x.stride(),
52	"dtype": x.dtype,	52	"dtype": x.dtype,
53	"device": x.device,	53	"device": x.device,
54	"storage_offset": x.storage_offset(), }	54	"storage_offset": x.storage_offset(), }
55		55
56	def reconstruct_from_tensor_metadata(metadata):	56	def reconstruct_from_tensor_metadata(metadata):
57	s = torch._C._construct_storage_from_data_pointer(	57	s = torch._C._construct_storage_from_data_pointer(
58	metadata["data_ptr"], metadata["device"], metadata["nbytes"])	58	metadata["data_ptr"], metadata["device"], metadata["nbytes"])
59	t = torch.empty([0], device=metadata["device"], dtype=metadata["dtype"])	59	t = torch.empty([0], device=metadata["device"], dtype=metadata["dtype"])
60	t.set_(source=s, storage_offset=metadata["storage_offset"],	60	t.set_(source=s, storage_offset=metadata["storage_offset"],
61	size=metadata["size"], stride=metadata["stride"], )	61	size=metadata["size"], stride=metadata["stride"], )
62	return t	62	return t
63		63
64	def cudagraphify(fn, inputs, pool, stream):	64	def cudagraphify(fn, inputs, pool, stream):
65	torch.npu.synchronize()	65	torch.npu.synchronize()
66	gc.collect()	66	gc.collect()
67	torch.npu.empty_cache()	67	torch.npu.empty_cache()
68		68
69	graph = torch.npu.NPUGraph()	69	graph = torch.npu.NPUGraph()
70	with torch.npu.graph(graph, stream=stream, pool=pool):	70	with torch.npu.graph(graph, stream=stream, pool=pool):
71	static_outputs = fn(*inputs)	71	static_outputs = fn(*inputs)
72	return graph, static_outputs	72	return graph, static_outputs
73		73
74	def foo(x, idx):	74	def foo(x, idx):
75	r1 = x.expand([1, 2097152 // 8]).sqrt()	75	r1 = x.expand([1, 2097152 // 8]).sqrt()
76	r2 = x.expand([idx, 2097152]).clone()	76	r2 = x.expand([idx, 2097152]).clone()
77	return r1, r2	77	return r1, r2
78		78
79	# init	79	# init
80	pool_id = torch.npu.graph_pool_handle()	80	pool_id = torch.npu.graph_pool_handle()
81	com_stream = torch.npu.Stream()	81	com_stream = torch.npu.Stream()
82	com_device = torch_npu.npu.current_device()	82	com_device = torch_npu.npu.current_device()
83	inp = torch.tensor([7]).npu()	83	inp = torch.tensor([7]).npu()
84		84
85	# start capture graph1	85	# start capture graph1
86	graph1, outputs1 = cudagraphify(foo, [inp, 1], pool=pool_id, stream=com_stream)	86	graph1, outputs1 = cudagraphify(foo, [inp, 1], pool=pool_id, stream=com_stream)
87	graph1_state = torch_npu._C._npu_getCheckpointState(com_device, pool_id)	87	graph1_state = torch_npu._C._npu_getCheckpointState(com_device, pool_id)
88	output1_metadata = [tensor_metadata(t) for t in outputs1]	88	output1_metadata = [tensor_metadata(t) for t in outputs1]
89	outputs1 = None	89	outputs1 = None
90		90
91	# start capture graph2	91	# start capture graph2
92	graph2, outputs2 = cudagraphify(foo, [inp, 2], pool=pool_id, stream=com_stream)	92	graph2, outputs2 = cudagraphify(foo, [inp, 2], pool=pool_id, stream=com_stream)
93	graph2_state = torch_npu._C._npu_getCheckpointState(com_device, pool_id)	93	graph2_state = torch_npu._C._npu_getCheckpointState(com_device, pool_id)
94	graph2.replay()	94	graph2.replay()
95	outputs2 = None	95	outputs2 = None
96		96
97	# replay graph1	97	# replay graph1
98	graph1.replay()	98	graph1.replay()
99	reconstructed_tensors1 = [reconstruct_from_tensor_metadata(metadata) for metadata in output1_metadata]	99	reconstructed_tensors1 = [reconstruct_from_tensor_metadata(metadata) for metadata in output1_metadata]
100	output1_new_storage = [output.untyped_storage()._cdata for output in reconstructed_tensors1]	100	output1_new_storage = [output.untyped_storage()._cdata for output in reconstructed_tensors1]
101	torch_npu._C._npu_setCheckpointPoolState(com_device, graph1_state, [], output1_new_storage)	101	torch_npu._C._npu_setCheckpointPoolState(com_device, graph1_state, [], output1_new_storage)
102		102
103		103
104	if __name__ == '__main__':	104	if __name__ == '__main__':
105	run_tests()	105	run_tests()

Mtest/npu/test_kernel_check.py+0-1

	@@ -118,4 +118,3 @@ class TestKernelCheck(TestCase):
118		118
119	if __name__ == "__main__":	119	if __name__ == "__main__":
120	run_tests()	120	run_tests()
121

Mtest/npu/test_pin_memory_host_register.py+19-19

Mtest/npu/test_save_async.py+119-119

	@@ -1,119 +1,119 @@
1	import os	1	import os
2	import time	2	import time
3	import copy	3	import copy
4		4
5	import torch	5	import torch
6	import torch.nn as nn	6	import torch.nn as nn
7	import torch.optim as optim	7	import torch.optim as optim
8		8
9	import torch_npu	9	import torch_npu
10	from torch_npu.testing.testcase import TestCase, run_tests	10	from torch_npu.testing.testcase import TestCase, run_tests
11	from torch_npu.utils._path_manager import PathManager	11	from torch_npu.utils._path_manager import PathManager
12		12
13		13
14	class TestAsyncSave(TestCase):	14	class TestAsyncSave(TestCase):
15	test_save_path = os.path.join(	15	test_save_path = os.path.join(
16	os.path.realpath(os.path.dirname(__file__)), "test_save_async")	16	os.path.realpath(os.path.dirname(__file__)), "test_save_async")
17		17
18	@classmethod	18	@classmethod
19	def setUpClass(cls):	19	def setUpClass(cls):
20	PathManager.make_dir_safety(TestAsyncSave.test_save_path)	20	PathManager.make_dir_safety(TestAsyncSave.test_save_path)
21		21
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
29	while time.time() - start_time < timeout_sec:	29	while time.time() - start_time < timeout_sec:
30	if os.path.exists(file_path):	30	if os.path.exists(file_path):
31	current_size = os.path.getsize(file_path)	31	current_size = os.path.getsize(file_path)
32	time.sleep(poll_interval_sec)	32	time.sleep(poll_interval_sec)
33	new_size = os.path.getsize(file_path)	33	new_size = os.path.getsize(file_path)
34		34
35	if current_size == new_size:	35	if current_size == new_size:
36	return True	36	return True
37	else:	37	else:
38	time.sleep(poll_interval_sec)	38	time.sleep(poll_interval_sec)
39		39
40	return False	40	return False
41		41
42	def test_save_async_tensor(self):	42	def test_save_async_tensor(self):
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 = []
56	model_list = []	56	model_list = []
57	checkpoint_list = []	57	checkpoint_list = []
58	model_origin = nn.Sequential(	58	model_origin = nn.Sequential(
59	nn.Linear(100, 50),	59	nn.Linear(100, 50),
60	nn.ReLU(),	60	nn.ReLU(),
61	nn.Linear(50, 20),	61	nn.Linear(50, 20),
62	nn.ReLU(),	62	nn.ReLU(),
63	nn.Linear(20, 5),	63	nn.Linear(20, 5),
64	nn.ReLU()	64	nn.ReLU()
65	)	65	)
66		66
67	input_data = torch.ones(6400, 100).npu()	67	input_data = torch.ones(6400, 100).npu()
68	labels = torch.arange(5).repeat(1280).npu()	68	labels = torch.arange(5).repeat(1280).npu()
69		69
70	criterion = nn.CrossEntropyLoss()	70	criterion = nn.CrossEntropyLoss()
71	model = model_origin.npu()	71	model = model_origin.npu()
72	optimerizer = optim.SGD(model.parameters(), lr=0.1)	72	optimerizer = optim.SGD(model.parameters(), lr=0.1)
73	for step in range(3):	73	for step in range(3):
74	outputs = model(input_data)	74	outputs = model(input_data)
75	loss = criterion(outputs, labels)	75	loss = criterion(outputs, labels)
76		76
77	optimerizer.zero_grad()	77	optimerizer.zero_grad()
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(),
85	"optimizer": optimerizer.state_dict()	85	"optimizer": optimerizer.state_dict()
86	}	86	}
87	checkpoint_list.append(copy.deepcopy(checkpoint))	87	checkpoint_list.append(copy.deepcopy(checkpoint))
88	model_list.append(copy.deepcopy(model))	88	model_list.append(copy.deepcopy(model))
89	checkpoint_async_path = os.path.join(TestAsyncSave.test_save_path, f"checkpoint_async_{step}.path")	89	checkpoint_async_path = os.path.join(TestAsyncSave.test_save_path, f"checkpoint_async_{step}.path")
90	model_async_path = os.path.join(TestAsyncSave.test_save_path, f"model_async_{step}.path")	90	model_async_path = os.path.join(TestAsyncSave.test_save_path, f"model_async_{step}.path")
91	torch_npu.utils.save_async(checkpoint, checkpoint_async_path, model=model)	91	torch_npu.utils.save_async(checkpoint, checkpoint_async_path, model=model)
92	torch_npu.utils.save_async(model, model_async_path, model=model)	92	torch_npu.utils.save_async(model, model_async_path, model=model)
93		93
94	for i in range(3):	94	for i in range(3):
95	self.assertEqual(loss1[i], loss2[i].item())	95	self.assertEqual(loss1[i], loss2[i].item())
96	checkpoint_async_path = os.path.join(TestAsyncSave.test_save_path, f"checkpoint_async_{i}.path")	96	checkpoint_async_path = os.path.join(TestAsyncSave.test_save_path, f"checkpoint_async_{i}.path")
97	if self.wait_for_save_completion(checkpoint_async_path):	97	if self.wait_for_save_completion(checkpoint_async_path):
98	checkpoint_async = torch.load(checkpoint_async_path, weights_only=False)	98	checkpoint_async = torch.load(checkpoint_async_path, weights_only=False)
99	self.assertEqual(checkpoint_list[i], checkpoint_async, prec=2e-3)	99	self.assertEqual(checkpoint_list[i], checkpoint_async, prec=2e-3)
100	else:	100	else:
101	self.assertTrue(False, f"{checkpoint_async_path} is not exist!")	101	self.assertTrue(False, f"{checkpoint_async_path} is not exist!")
102	model_async_path = os.path.join(TestAsyncSave.test_save_path, f"model_async_{i}.path")	102	model_async_path = os.path.join(TestAsyncSave.test_save_path, f"model_async_{i}.path")
103	if self.wait_for_save_completion(model_async_path):	103	if self.wait_for_save_completion(model_async_path):
104	model_async = torch.load(model_async_path, weights_only=False)	104	model_async = torch.load(model_async_path, weights_only=False)
105	else:	105	else:
106	self.assertTrue(False, f"{model_async_path} is not exist!")	106	self.assertTrue(False, f"{model_async_path} is not exist!")
107	state_dict_sync = model_list[i].state_dict()	107	state_dict_sync = model_list[i].state_dict()
108	state_dict_async = model_async.state_dict()	108	state_dict_async = model_async.state_dict()
109		109
110	key_sync = sorted(state_dict_sync.keys())	110	key_sync = sorted(state_dict_sync.keys())
111	key_async = sorted(state_dict_async.keys())	111	key_async = sorted(state_dict_async.keys())
112		112
113	self.assertEqual(key_sync, key_async)	113	self.assertEqual(key_sync, key_async)
114	for key in key_async:	114	for key in key_async:
115	self.assertEqual(state_dict_async[key], state_dict_sync[key], prec=2e-3)	115	self.assertEqual(state_dict_async[key], state_dict_sync[key], prec=2e-3)
116		116
117	if __name__ == '__main__':	117	if __name__ == '__main__':
118	torch.npu.set_device(0)	118	torch.npu.set_device(0)
119	run_tests()	119	run_tests()

Mtest/profiler/analysis/test_profiler_config.py+0-1

	@@ -81,4 +81,3 @@ class TestProfilerConfig(TestCase):
81		81
82	if __name__ == "__main__":	82	if __name__ == "__main__":
83	run_tests()	83	run_tests()
84

Mtest/profiler/test_non_intrusive_profile.py+0-1

	@@ -39,4 +39,3 @@ class TestNoneInstrusiveProfile(TestCase):
39		39
40	if __name__ == "__main__":	40	if __name__ == "__main__":
41	run_tests()	41	run_tests()
42

Mtest/test_fake_tensor.py+2016-2016

Mtest/unsupported_test_cases/.pytorch-disabled-tests.json+3-1

	@@ -32154,5 +32154,7 @@
32154	"test_to_complex_npu (__main__.TestNNDeviceTypePRIVATEUSE1)": ["", [""]],	32154	"test_to_complex_npu (__main__.TestNNDeviceTypePRIVATEUSE1)": ["", [""]],
32155	"test_npu_roi_align_1 (__main__.TestPsRoiPooling)": ["", [""]],	32155	"test_npu_roi_align_1 (__main__.TestPsRoiPooling)": ["", [""]],
32156	"test_silu (__main__.TestActivations)": ["", [""]],	32156	"test_silu (__main__.TestActivations)": ["", [""]],
32157	"test_codegen_upcast_to_fp32_emits_cast_bfloat16_upcast_flag_True (__main__.TestCodegenUpcastToFP32)": ["", [""]]	32157	"test_codegen_upcast_to_fp32_emits_cast_bfloat16_upcast_flag_True (__main__.TestCodegenUpcastToFP32)": ["", [""]],
		32158	"test_data_parallel_rnn (__main__.TestDataParallel)": ["", ["Disabled during A1 to A2 chip transition"]],
		32159	"test_alltoall_single_2p_size_dist (__main__.HcclAlltoAllSingleTest)": ["", ["Disabled during A1 to A2 chip transition"]]
32158	}	32160	}

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

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

	@@ -1,330 +1,330 @@
1	import os	1	import os
2	import sys	2	import sys
3	import importlib	3	import importlib
4	import shutil	4	import shutil
5	from itertools import count	5	from itertools import count
6	from typing import Any, Callable, Dict, List, Optional, Tuple, Iterator	6	from typing import Any, Callable, Dict, List, Optional, Tuple, Iterator
7		7
8	import torch	8	import torch
9	from torch._inductor.compile_fx import clone_preserve_strides	9	from torch._inductor.compile_fx import clone_preserve_strides
10		10
11	from .. import config as anir_config	11	from .. import config as anir_config
12	from .utils import replace_placeholders	12	from .utils import replace_placeholders
13		13
14		14
15	_dump_id_iter: Iterator[int] = count()	15	_dump_id_iter: Iterator[int] = count()
16		16
17	class MetaCompiler:	17	class MetaCompiler:
18	def __init__(	18	def __init__(
19	self,	19	self,
20	kernel_name: str = "",	20	kernel_name: str = "",
21	multiprocess_compile: bool = False,	21	multiprocess_compile: bool = False,
22	no_more_compile: bool = False,	22	no_more_compile: bool = False,
23	kernel_meta: Optional[Dict[str, Any]] = None,	23	kernel_meta: Optional[Dict[str, Any]] = None,
24	autotune: bool = True,	24	autotune: bool = True,
25	):	25	):
26	kernel_meta = kernel_meta or {}	26	kernel_meta = kernel_meta or {}
27		27
28	self.kernel_name = kernel_name	28	self.kernel_name = kernel_name
29	self.kernel_meta = kernel_meta	29	self.kernel_meta = kernel_meta
30		30
31	self.dynamic = kernel_meta.get("dynamic")	31	self.dynamic = kernel_meta.get("dynamic")
32	self.mutated_indices = kernel_meta.get("mutated_indices") or []	32	self.mutated_indices = kernel_meta.get("mutated_indices") or []
33	self.kernel_hash = kernel_meta.get("kernel_hash")	33	self.kernel_hash = kernel_meta.get("kernel_hash")
34	self.signature = kernel_meta.get("signature")	34	self.signature = kernel_meta.get("signature")
35	self.ranks = kernel_meta.get("ranks")	35	self.ranks = kernel_meta.get("ranks")
36	self.num_outputs = kernel_meta.get("num_outputs")	36	self.num_outputs = kernel_meta.get("num_outputs")
37	self.num_call_functions = kernel_meta.get("num_call_functions")	37	self.num_call_functions = kernel_meta.get("num_call_functions")
38	self.device_index = kernel_meta.get("device_index", 0)	38	self.device_index = kernel_meta.get("device_index", 0)
39	self.traced_graph_hash = kernel_meta.get("traced_graph_hash")	39	self.traced_graph_hash = kernel_meta.get("traced_graph_hash")
40		40
41	self.multiprocess_compile = multiprocess_compile	41	self.multiprocess_compile = multiprocess_compile
42	self.no_more_compile = no_more_compile	42	self.no_more_compile = no_more_compile
43	self.autotune = autotune	43	self.autotune = autotune
44		44
45	self._fallback_call: Optional[Callable[..., Any]] = None	45	self._fallback_call: Optional[Callable[..., Any]] = None
46	self._nc_input_indices: Optional[List[int]] = None	46	self._nc_input_indices: Optional[List[int]] = None
47	self._nc_output_indices: Optional[List[int]] = None	47	self._nc_output_indices: Optional[List[int]] = None
48		48
49	self.autotuned: bool = False	49	self.autotuned: bool = False
50	self.launchers: list[Callable[..., Any]] = []	50	self.launchers: list[Callable[..., Any]] = []
51	self.kernel_paths: list[Any] = []	51	self.kernel_paths: list[Any] = []
52	self.is_fallback_kernels: list[bool] = []	52	self.is_fallback_kernels: list[bool] = []
53		53
54	def register_launcher(	54	def register_launcher(
55	self,	55	self,
56	launcher: Callable[..., Any],	56	launcher: Callable[..., Any],
57	kernel_path: Any = None,	57	kernel_path: Any = None,
58	is_fallback_kernel: bool = False,	58	is_fallback_kernel: bool = False,
59	) -> None:	59	) -> None:
60	"""Register a compiled or fallback launcher for runtime execution."""	60	"""Register a compiled or fallback launcher for runtime execution."""
61	self.launchers.append(launcher)	61	self.launchers.append(launcher)
62	self.kernel_paths.append(kernel_path)	62	self.kernel_paths.append(kernel_path)
63	self.is_fallback_kernels.append(is_fallback_kernel)	63	self.is_fallback_kernels.append(is_fallback_kernel)
64		64
65	def get_primary_launcher_index(self) -> int:	65	def get_primary_launcher_index(self) -> int:
66	if not self.launchers:	66	if not self.launchers:
67	raise RuntimeError("No valid launcher")	67	raise RuntimeError("No valid launcher")
68	for idx, is_fallback_kernel in enumerate(self.is_fallback_kernels):	68	for idx, is_fallback_kernel in enumerate(self.is_fallback_kernels):
69	if not is_fallback_kernel:	69	if not is_fallback_kernel:
70	return idx	70	return idx
71	return 0	71	return 0
72		72
73	def get_primary_launcher(self) -> Callable[..., Any]:	73	def get_primary_launcher(self) -> Callable[..., Any]:
74	return self.launchers[self.get_primary_launcher_index()]	74	return self.launchers[self.get_primary_launcher_index()]
75		75
76	def data_dump(self, *args, dump_path=None):	76	def data_dump(self, *args, dump_path=None):
77	if not dump_path:	77	if not dump_path:
78	dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(self.device_index), self.kernel_name)	78	dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(self.device_index), self.kernel_name)
79	data_dump_path = os.path.join(dump_path, 'data.pth')	79	data_dump_path = os.path.join(dump_path, 'data.pth')
80	args_cpu = [arg.cpu() if isinstance(arg, torch.Tensor) else arg for arg in args]	80	args_cpu = [arg.cpu() if isinstance(arg, torch.Tensor) else arg for arg in args]
81	torch.save(args_cpu, data_dump_path)	81	torch.save(args_cpu, data_dump_path)
82		82
83	def data_dump_fake(self, *args, dump_path=None):	83	def data_dump_fake(self, *args, dump_path=None):
84	if not dump_path:	84	if not dump_path:
85	dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(self.device_index), self.kernel_name)	85	dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(self.device_index), self.kernel_name)
86	runable_py_path = os.path.join(dump_path, f'runnable_{self.kernel_name}.py')	86	runable_py_path = os.path.join(dump_path, f'runnable_{self.kernel_name}.py')
87	fake_inputs = [f'rand_strided({arg.shape}, {arg.stride()}, device="{arg.device.type}", dtype={arg.dtype})' \	87	fake_inputs = [f'rand_strided({arg.shape}, {arg.stride()}, device="{arg.device.type}", dtype={arg.dtype})' \
88	if isinstance(arg, torch.Tensor) else str(arg) for arg in args[:-self.num_outputs]]	88	if isinstance(arg, torch.Tensor) else str(arg) for arg in args[:-self.num_outputs]]
89	fake_outputs = [f'empty_strided({arg.shape}, {arg.stride()}, device="{arg.device.type}", dtype={arg.dtype})' \	89	fake_outputs = [f'empty_strided({arg.shape}, {arg.stride()}, device="{arg.device.type}", dtype={arg.dtype})' \
90	if isinstance(arg, torch.Tensor) else str(arg) for arg in args[-self.num_outputs:]]	90	if isinstance(arg, torch.Tensor) else str(arg) for arg in args[-self.num_outputs:]]
91	replacements = {"FAKE_ARGS_PLACEHOLDER": f"args = [{', '.join(fake_inputs + fake_outputs)}]"}	91	replacements = {"FAKE_ARGS_PLACEHOLDER": f"args = [{', '.join(fake_inputs + fake_outputs)}]"}
92	replace_placeholders(runable_py_path, replacements)	92	replace_placeholders(runable_py_path, replacements)
93		93
94	def fx_subgraph_dump(self, suffix):	94	def fx_subgraph_dump(self, suffix):
95	subgraph_dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(self.device_index), self.kernel_name)	95	subgraph_dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(self.device_index), self.kernel_name)
96	failed_fx_subgraph_dump_path = anir_config.fx_subgraph_dump_path + f'_{suffix}'	96	failed_fx_subgraph_dump_path = anir_config.fx_subgraph_dump_path + f'_{suffix}'
97	failed_subgraph_dump_path = os.path.join(failed_fx_subgraph_dump_path, str(self.device_index), f'{next(_dump_id_iter)}_' + self.kernel_name)	97	failed_subgraph_dump_path = os.path.join(failed_fx_subgraph_dump_path, str(self.device_index), f'{next(_dump_id_iter)}_' + self.kernel_name)
98	if os.path.exists(failed_subgraph_dump_path):	98	if os.path.exists(failed_subgraph_dump_path):
99	shutil.rmtree(failed_subgraph_dump_path)	99	shutil.rmtree(failed_subgraph_dump_path)
100	shutil.copytree(subgraph_dump_path, failed_subgraph_dump_path)	100	shutil.copytree(subgraph_dump_path, failed_subgraph_dump_path)
101	return failed_subgraph_dump_path	101	return failed_subgraph_dump_path
102		102
103	def acc_compare_and_dump(self, args, *kwargs):	103	def acc_compare_and_dump(self, args, *kwargs):
104	from torch.testing._comparison import _make_mismatch_msg	104	from torch.testing._comparison import _make_mismatch_msg
105	self.register_fx_fallback(self.kernel_meta)	105	self.register_fx_fallback(self.kernel_meta)
106	launcher_fx = self.launchers[1]	106	launcher_fx = self.launchers[1]
107	launcher = self.launchers[0]	107	launcher = self.launchers[0]
108		108
109	fx_outputs = [clone_preserve_strides(arg).to(torch.float32) if arg.dtype == torch.bfloat16 \	109	fx_outputs = [clone_preserve_strides(arg).to(torch.float32) if arg.dtype == torch.bfloat16 \
110	else clone_preserve_strides(arg) for arg in args[-self.num_outputs:]]	110	else clone_preserve_strides(arg) for arg in args[-self.num_outputs:]]
111	fx_inputs = [clone_preserve_strides(arg) if isinstance(arg, torch.Tensor) else arg for arg in args[:-self.num_outputs]]	111	fx_inputs = [clone_preserve_strides(arg) if isinstance(arg, torch.Tensor) else arg for arg in args[:-self.num_outputs]]
112	fx_inputs = [inp.float() if isinstance(inp, torch.Tensor) and inp.dtype == torch.bfloat16 else inp for inp in fx_inputs]	112	fx_inputs = [inp.float() if isinstance(inp, torch.Tensor) and inp.dtype == torch.bfloat16 else inp for inp in fx_inputs]
113		113
114	fx_args = fx_inputs + fx_outputs	114	fx_args = fx_inputs + fx_outputs
115	launcher_fx(fx_args, *kwargs)	115	launcher_fx(fx_args, *kwargs)
116		116
117	if self.dynamic:	117	if self.dynamic:
118	args_new = self.prepare_runtime_args(	118	args_new = self.prepare_runtime_args(
119	list(args),	119	list(args),
120	)	120	)
121	else:	121	else:
122	args_new = args	122	args_new = args
123		123
124	output = launcher(args_new, *kwargs)	124	output = launcher(args_new, *kwargs)
125		125
126	has_acc_error = False	126	has_acc_error = False
127	num_inputs = len(args) - self.num_outputs	127	num_inputs = len(args) - self.num_outputs
128	for idx, (actual, expected) in enumerate(zip(args[num_inputs:], fx_outputs)):	128	for idx, (actual, expected) in enumerate(zip(args[num_inputs:], fx_outputs)):
129	if actual.dtype != expected.dtype:	129	if actual.dtype != expected.dtype:
130	expected = expected.to(actual.dtype)	130	expected = expected.to(actual.dtype)
131	acc_comp_tol = anir_config.acc_comp_tol.get(actual.dtype, anir_config.acc_comp_tol['default'])	131	acc_comp_tol = anir_config.acc_comp_tol.get(actual.dtype, anir_config.acc_comp_tol['default'])
132	rtol = acc_comp_tol['rtol']	132	rtol = acc_comp_tol['rtol']
133	atol = acc_comp_tol['atol']	133	atol = acc_comp_tol['atol']
134	matches = torch.isclose(	134	matches = torch.isclose(
135	actual, expected, rtol=rtol, atol=atol, equal_nan=True	135	actual, expected, rtol=rtol, atol=atol, equal_nan=True
136	)	136	)
137	if not matches.all():	137	if not matches.all():
138	abs_diff = abs(actual - expected)	138	abs_diff = abs(actual - expected)
139	rel_diff = abs_diff / abs(expected)	139	rel_diff = abs_diff / abs(expected)
140	rel_diff.masked_fill_(matches, 0)	140	rel_diff.masked_fill_(matches, 0)
141	number_of_elements = matches.numel()	141	number_of_elements = matches.numel()
142	total_mismatches = number_of_elements - int(torch.sum(matches))	142	total_mismatches = number_of_elements - int(torch.sum(matches))
143	extra = (	143	extra = (
144	f"Mismatched elements: {total_mismatches} / {number_of_elements} "	144	f"Mismatched elements: {total_mismatches} / {number_of_elements} "
145	f"({total_mismatches / number_of_elements:.1%})"	145	f"({total_mismatches / number_of_elements:.1%})"
146	)	146	)
147	msg = _make_mismatch_msg(	147	msg = _make_mismatch_msg(
148	default_identifier="Tensor-likes",	148	default_identifier="Tensor-likes",
149	identifier=None,	149	identifier=None,
150	extra=extra,	150	extra=extra,
151	abs_diff=abs_diff.max().item(),	151	abs_diff=abs_diff.max().item(),
152	abs_diff_idx=None,	152	abs_diff_idx=None,
153	atol=atol,	153	atol=atol,
154	rel_diff=rel_diff.max().item(),	154	rel_diff=rel_diff.max().item(),
155	rel_diff_idx=None,	155	rel_diff_idx=None,
156	rtol=rtol,	156	rtol=rtol,
157	)	157	)
158	print(f"Kernel Name: {self.kernel_name}\n{msg}", flush=True)	158	print(f"Kernel Name: {self.kernel_name}\n{msg}", flush=True)
159	has_acc_error = True	159	has_acc_error = True
160		160
161	del abs_diff	161	del abs_diff
162	del rel_diff	162	del rel_diff
163	del matches	163	del matches
164	del expected	164	del expected
165		165
166	if anir_config.fx_subgraph_dump_path:	166	if anir_config.fx_subgraph_dump_path:
167	data = args	167	data = args
168	if has_acc_error:	168	if has_acc_error:
169	data_dump_path = self.fx_subgraph_dump('acc_failed')	169	data_dump_path = self.fx_subgraph_dump('acc_failed')
170	self.data_dump_fake(*data, dump_path=data_dump_path)	170	self.data_dump_fake(*data, dump_path=data_dump_path)
171	del fx_inputs	171	del fx_inputs
172	torch.npu.synchronize()	172	torch.npu.synchronize()
173	self.launchers = [self.launchers[0]]	173	self.launchers = [self.launchers[0]]
174	self.is_fallback_kernels = [self.is_fallback_kernels[0]]	174	self.is_fallback_kernels = [self.is_fallback_kernels[0]]
175		175
176	return output	176	return output
177		177
178	def compile(self, args, *kwargs):	178	def compile(self, args, *kwargs):
179	raise NotImplementedError	179	raise NotImplementedError
180		180
181	def ensure_runtime_ready(self, args, *kwargs) -> None:	181	def ensure_runtime_ready(self, args, *kwargs) -> None:
182	if not self.autotuned:	182	if not self.autotuned:
183	if self.autotune:	183	if self.autotune:
184	self.register_fx_fallback(self.kernel_meta)	184	self.register_fx_fallback(self.kernel_meta)
185	self.autotune_to_one_config(args, *kwargs)	185	self.autotune_to_one_config(args, *kwargs)
186	self.autotuned = True	186	self.autotuned = True
187		187
188	def prepare_runtime_args(	188	def prepare_runtime_args(
189	self,	189	self,
190	args_list: List[Any],	190	args_list: List[Any],
191	) -> List[Any]:	191	) -> List[Any]:
192	args_new = ()	192	args_new = ()
193	for arg in args_list:	193	for arg in args_list:
194	if not torch.is_tensor(arg):	194	if not torch.is_tensor(arg):
195	args_new = args_new + (arg,)	195	args_new = args_new + (arg,)
196	continue	196	continue
197	args_new = args_new + (arg, arg, 0) + arg.size() + arg.stride()	197	args_new = args_new + (arg, arg, 0) + arg.size() + arg.stride()
198	args_list = list(args_new)	198	args_list = list(args_new)
199	return args_list	199	return args_list
200		200
201	def autotune_to_one_config(self, args, *kwargs):	201	def autotune_to_one_config(self, args, *kwargs):
202	return None	202	return None
203		203
204	def _normalize_contiguous_args(	204	def _normalize_contiguous_args(
205	self, args: List[Any]	205	self, args: List[Any]
206	) -> Tuple[List[Any], Optional[List[torch.Tensor]]]:	206	) -> Tuple[List[Any], Optional[List[torch.Tensor]]]:
207	meta = self.kernel_meta or {}	207	meta = self.kernel_meta or {}
208	num_outputs = int(meta.get("num_outputs") or 0)	208	num_outputs = int(meta.get("num_outputs") or 0)
209	num_call_functions = int(meta.get("num_call_functions") or 0)	209	num_call_functions = int(meta.get("num_call_functions") or 0)
210		210
211	if num_outputs < 0 or len(args) < num_outputs:	211	if num_outputs < 0 or len(args) < num_outputs:
212	return args, None	212	return args, None
213		213
214	num_inputs = len(args) - num_outputs	214	num_inputs = len(args) - num_outputs
215		215
216	if self._nc_input_indices is None:	216	if self._nc_input_indices is None:
217	self._nc_input_indices = []	217	self._nc_input_indices = []
218	if num_call_functions > 0:	218	if num_call_functions > 0:
219	input_slice = args[:num_inputs] if num_outputs > 0 else args	219	input_slice = args[:num_inputs] if num_outputs > 0 else args
220	for idx, arg in enumerate(input_slice):	220	for idx, arg in enumerate(input_slice):
221	if not isinstance(arg, torch.Tensor) or arg.is_contiguous():	221	if not isinstance(arg, torch.Tensor) or arg.is_contiguous():
222	continue	222	continue
223	args[idx] = arg.contiguous()	223	args[idx] = arg.contiguous()
224	self._nc_input_indices.append(idx)	224	self._nc_input_indices.append(idx)
225	else:	225	else:
226	for idx in self._nc_input_indices:	226	for idx in self._nc_input_indices:
227	if idx < len(args) and isinstance(args[idx], torch.Tensor):	227	if idx < len(args) and isinstance(args[idx], torch.Tensor):
228	args[idx] = args[idx].contiguous()	228	args[idx] = args[idx].contiguous()
229		229
230	if num_outputs == 0:	230	if num_outputs == 0:
231	return args, None	231	return args, None
232		232
233	original_outputs: Optional[List[torch.Tensor]] = None	233	original_outputs: Optional[List[torch.Tensor]] = None
234	if self._nc_output_indices is None:	234	if self._nc_output_indices is None:
235	self._nc_output_indices = []	235	self._nc_output_indices = []
236	original_outputs = []	236	original_outputs = []
237	for j, out in enumerate(args[num_inputs:]):	237	for j, out in enumerate(args[num_inputs:]):
238	if not isinstance(out, torch.Tensor) or out.is_contiguous():	238	if not isinstance(out, torch.Tensor) or out.is_contiguous():
239	continue	239	continue
240	tmp = torch.empty(out.shape, dtype=out.dtype, device=out.device)	240	tmp = torch.empty(out.shape, dtype=out.dtype, device=out.device)
241	out_idx = num_inputs + j	241	out_idx = num_inputs + j
242	original_outputs.append(out)	242	original_outputs.append(out)
243	args[out_idx] = tmp	243	args[out_idx] = tmp
244	self._nc_output_indices.append(out_idx)	244	self._nc_output_indices.append(out_idx)
245	else:	245	else:
246	original_outputs = []	246	original_outputs = []
247	for out_idx in self._nc_output_indices:	247	for out_idx in self._nc_output_indices:
248	tmp = torch.empty(	248	tmp = torch.empty(
249	args[out_idx].shape, dtype=args[out_idx].dtype, device=args[out_idx].device	249	args[out_idx].shape, dtype=args[out_idx].dtype, device=args[out_idx].device
250	)	250	)
251	original_outputs.append(args[out_idx])	251	original_outputs.append(args[out_idx])
252	args[out_idx] = tmp	252	args[out_idx] = tmp
253		253
254	return args, original_outputs	254	return args, original_outputs
255		255
256	def _copy_back_non_contiguous_outputs(	256	def _copy_back_non_contiguous_outputs(
257	self,	257	self,
258	args_list: List[Any],	258	args_list: List[Any],
259	original_outputs: Optional[List[torch.Tensor]],	259	original_outputs: Optional[List[torch.Tensor]],
260	) -> None:	260	) -> None:
261	if original_outputs and self._nc_output_indices:	261	if original_outputs and self._nc_output_indices:
262	for orig, idx in zip(original_outputs, self._nc_output_indices):	262	for orig, idx in zip(original_outputs, self._nc_output_indices):
263	orig.copy_(args_list[idx])	263	orig.copy_(args_list[idx])
264		264
265		265
266	def _fx_graph_call_factory(self, module: torch.nn.Module, num_outputs: int) -> Callable[..., Any]:	266	def _fx_graph_call_factory(self, module: torch.nn.Module, num_outputs: int) -> Callable[..., Any]:
267	def module_call(args, *kwargs):	267	def module_call(args, *kwargs):
268	actual_args = args[:-num_outputs] if num_outputs > 0 else args	268	actual_args = args[:-num_outputs] if num_outputs > 0 else args
269	actual_outputs = module.forward(*actual_args)	269	actual_outputs = module.forward(*actual_args)
270	for out1, out2 in zip(actual_outputs, args[-num_outputs:]):	270	for out1, out2 in zip(actual_outputs, args[-num_outputs:]):
271	if isinstance(out1, torch.Tensor) and not out1.is_contiguous():	271	if isinstance(out1, torch.Tensor) and not out1.is_contiguous():
272	out1 = out1.contiguous()	272	out1 = out1.contiguous()
273	out2.data = out1.data	273	out2.data = out1.data
274	return module_call	274	return module_call
275		275
276	def _load_traced_graph_model(self) -> torch.nn.Module:	276	def _load_traced_graph_model(self) -> torch.nn.Module:
277	meta = self.kernel_meta or {}	277	meta = self.kernel_meta or {}
278	traced_graph_hash = meta.get("traced_graph_hash")	278	traced_graph_hash = meta.get("traced_graph_hash")
279	if not traced_graph_hash:	279	if not traced_graph_hash:
280	raise RuntimeError("traced_graph_hash missing, cannot build FX fallback")	280	raise RuntimeError("traced_graph_hash missing, cannot build FX fallback")
281		281
282	traced_graph_cache = meta.get("traced_graph_cache")	282	traced_graph_cache = meta.get("traced_graph_cache")
283	if traced_graph_cache is None:	283	if traced_graph_cache is None:
284	raise RuntimeError("traced_graph_cache missing, cannot locate traced graph dump")	284	raise RuntimeError("traced_graph_cache missing, cannot locate traced graph dump")
285		285
286	base_cache = os.getenv("TORCHINDUCTOR_CACHE_DIR", "")	286	base_cache = os.getenv("TORCHINDUCTOR_CACHE_DIR", "")
287	dump_root = os.path.join(base_cache, traced_graph_cache)	287	dump_root = os.path.join(base_cache, traced_graph_cache)
288	dump_path = os.path.join(dump_root, str(self.device_index), traced_graph_hash)	288	dump_path = os.path.join(dump_root, str(self.device_index), traced_graph_hash)
289		289
290	sys.path.append(dump_path)	290	sys.path.append(dump_path)
291	try:	291	try:
292	module = importlib.import_module(traced_graph_hash)	292	module = importlib.import_module(traced_graph_hash)
293	finally:	293	finally:
294	sys.path.remove(dump_path)	294	sys.path.remove(dump_path)
295		295
296	Model = getattr(module, traced_graph_hash, None)	296	Model = getattr(module, traced_graph_hash, None)
297	if Model is None:	297	if Model is None:
298	raise RuntimeError("Cannot find valid graph module class in traced graph dump")	298	raise RuntimeError("Cannot find valid graph module class in traced graph dump")
299		299
300	return Model()	300	return Model()
301		301
302	def register_fx_fallback(self, kernel_meta) -> None:	302	def register_fx_fallback(self, kernel_meta) -> None:
303	model = self._load_traced_graph_model()	303	model = self._load_traced_graph_model()
304	num_outputs = kernel_meta.get("num_outputs", 0)	304	num_outputs = kernel_meta.get("num_outputs", 0)
305	module_call = self._fx_graph_call_factory(model, num_outputs)	305	module_call = self._fx_graph_call_factory(model, num_outputs)
306	self.register_launcher(	306	self.register_launcher(
307	module_call,	307	module_call,
308	kernel_path=self.kernel_name + "_fx_fallback",	308	kernel_path=self.kernel_name + "_fx_fallback",
309	is_fallback_kernel=True,	309	is_fallback_kernel=True,
310	)	310	)
311		311
312	def run(self, args, *kwargs):	312	def run(self, args, *kwargs):
313	args_list = list(args)	313	args_list = list(args)
314	args_list, original_outputs = self._normalize_contiguous_args(args_list)	314	args_list, original_outputs = self._normalize_contiguous_args(args_list)
315	self.ensure_runtime_ready(args_list, *kwargs)	315	self.ensure_runtime_ready(args_list, *kwargs)
316		316
317	launcher_idx = self.get_primary_launcher_index()	317	launcher_idx = self.get_primary_launcher_index()
318	launcher = self.launchers[launcher_idx]	318	launcher = self.launchers[launcher_idx]
319	is_fallback_kernel = self.is_fallback_kernels[launcher_idx]	319	is_fallback_kernel = self.is_fallback_kernels[launcher_idx]
320	if anir_config.online_acc_comp and not is_fallback_kernel:	320	if anir_config.online_acc_comp and not is_fallback_kernel:
321	output = self.acc_compare_and_dump(args_list, *kwargs)	321	output = self.acc_compare_and_dump(args_list, *kwargs)
322	self._copy_back_non_contiguous_outputs(args_list, original_outputs)	322	self._copy_back_non_contiguous_outputs(args_list, original_outputs)
323	return output	323	return output
324	if self.dynamic and not is_fallback_kernel:	324	if self.dynamic and not is_fallback_kernel:
325	args_list = self.prepare_runtime_args(	325	args_list = self.prepare_runtime_args(
326	args_list,	326	args_list,
327	)	327	)
328	ret = launcher(tuple(args_list), *kwargs)	328	ret = launcher(tuple(args_list), *kwargs)
329	self._copy_back_non_contiguous_outputs(args_list, original_outputs)	329	self._copy_back_non_contiguous_outputs(args_list, original_outputs)
330	return ret	330	return ret

Mtorch_npu/_inductor/docs/feature/non_contiguous_accesses/arch.drawio.svg+597-597

	@@ -1,598 +1,598 @@
1	<svg host="65bd71144e" xmlns="http://www.w3.org/2000/svg" style="background: transparent; background-color: transparent;" xmlns:xlink="http://www.w3.org/1999/xlink" version="1.1" width="1111px" height="601px" viewBox="-0.5 -0.5 1111 601" content="<mxfile><diagram id="uiDLYcUB2_ILZagOPzox" name="Page-1">7Vxbc5s4FP41THd3Jh7ul0d8SZuZps0k6Wz75MEg29oC8mI5sffXrwTCIBAOtnGcZMxDgo4uCH0fR+ccSZa0QbT+nHiL+S0KQCipcrCWtKGkqoqlkL9UsMkEhmZlglkCA1amEDzA/wATyky6ggFYcgUxQiGGC17oozgGPuZkXpKgZ77YFIX8UxfeDNQED74X1qV/wwDPM6ltyIX8C4Czef5kRWY5kZcXZoLl3AvQc0mkjSRtkCCEs7toPQAhHbt8XLJ61w25244lIMZtKqhZhScvXLF3Y/3Cm/xlE7SKA0DLy5LWf55DDB4Wnk9znwm6RDbHUUhSCrn1QjiLyb1Png8SIngCCYZk5FyWMUEYo4hkTFGMr70IhpQDX0D4BGg5lsEgV1SWHqAQJWl3tOyij1phtMzK0Y4lgCS8Sdpr2pMpDMNStev0IvIlTtBvUMox04vkhN4EhHekTQyR8B2+VgpgRN++PuQMBVoNrEsiBsFngCKAkw0pkufKOUE2PNGfC3YpSi6cl6hl5hU9RunZtvECdXLDgBeTQKthDgLCd5ZECZ6jGYq9cFRI+zwrijJfER2UFIF/AMYbhiQFi2cKGa1k85PW7xl58hdrLk0M11xqw1IH8SYAU28VYp41ioA1jWAu0Srx2egYTOV4yQzkpexMRgduJ+IJCD0Mn3hFcgx4+gW8o8FzzgWeIVC/ZvqqdAA4WM1/VyjPuMoGwSUFFH2xTl89zyd3s/T/yJJcU+qTG1NyDMnVpJEtuSOpfy2NSHIoOTYt01fSMhbNcqw8S877Qd4g6wpr9bSzQwSDIGVoa6LsNwsEHrCnft4MewtFOCv4NphMO1Lvms2rd1UW6HdVoN71DrS7dVEQeymIfJotawjVPJeGsC/oHY+e0TV6aVU3SbxNqcACwRgvSy3fUUFJDZgOrwZ0Xe8ZvHlerWJrxu4q5CbrR8Gm7Qu1IphzIdjxBJPfLcEM/dQEywdMYOQsF158lJETIi8gz19ilABSbvjwtWS3ZK1/CLtlavvAb2W3TGyCqNyR3WJqPJ80MyfHq1guilZjDg57DHNyx2CvIEteDh8KXxUG9sWRdo2+ZAzrKoQf/xjFoAJeLjpOcfGwd4CsUbFINU2ErCMAVusC2Ga/J4BPxTebiSZCFTHx/N+z9Cu+8rPxopoCxhBDLxQqCxwtym7NZPscbcho9QeMFziR1P7mz+3Dk2p3SNVSJyHRInRQ16SW/Vepyy88sz1rQzDFb5ezeiNnO1FAeiUuZgkcJ13AUqMLltr7T1wUxqtsAqF0jBGbTAT+eXeOvkMd975OHX17JLmGwNFv8uY/4mR5Jie/Olmq9is6+XkI681z1ZKc/oWr5+aqYp6Rq2qHDoHcnsQPN7ekQTkNh5pS36E3riLZ9sdk1plcBr3CLEMVugzWqbilCLhVjW/EgUsXgQtLJ/CW8xRYhQeRR/y0wQewhvgnexK9T8MhPcsxWLqIiNDEppS4AwkkA0VZVoRMsuCKQ9RpObxC2tOk3SGWNFVtspEZL0Y8SqAbO+y0/QIj9TBGdQHVkJWeyreSBXFYxR0hEdKY1bMsZ3tZXMu64/Sc8qXzj8niQrXHHBA4yalb4vIARQsYElhG11J/QHTXI1F6phdR5ZP9/Xb34+rmPs+vB/fer3fc7Gl07h3LFVPO0PW6CrNP5Byr9ajHTRysfEze+oLmIWgaFWPH0gRTkmhG6gROvQanS8YVr+KPFLp6RTjNiqrXFZGFcbKvszl0xRuPnTlfMA5gAnw8ZoHPbZqFPxuM1zfo/OUzVTxZLniPrpWjt1MsiBu+pXcnzscjcz5cm7od1AsZpm7vDndkv2HhhuDiu+yvWSzjvL6L+YJqOY1fnG4CciX3OuPoxRfubvlMb8Un+VR8OiB+vZtPDEiaw1YlhNOVYEFk3+np+H5s2lL5/Rpd7dYL5S6ovFV6+UKMIwgYiojcic1VD25f1oG7MaaV6qKFUcf1VKvAWj0O7GfxjXGEAl4fLGGExygON9yHXrX8aCnSYFZOxjCE8UxKj4FM4aylpfReF2JfkTSmWvHA9Fe1k7R6jHc3bzCIFqGHQRvuFGUvtOmYNrbeijYiddMJbepxtV20CV5iS3BRMafjil4N8oi5Itr60QlX6kG7IgYrtzpa8XHMEfMVYa9sYBVYI4poGakTc8SqYe7S2gRd8sf4QIC+5j5DpQKoVV9JOVXkXav7vxcYD7T4TKMKo0Adi7RxJ0DW/b8LkAfaYMo5gdTrKrYGY/n0CRvE8sGRfJdEL98X8auU07BH4sDDJlnXKjscXj4fwrQOd8BUFqPS1XGQIzYZ6HVn6ubb8Mfg8fv9mNzc3I8Gj+Pb0e33+1/j2+9D0sbwEzV+x6lnFcH1p8uHeNBaQ8XAFW/A1072Jdb15+4PrxX3dbPO/dOcvavvANIrkSylugW8Yf/PId9M80KNYBmSrR3620WGwq3U5PRqXndsHR9Xdx2QKk55p7uCHNfDIO6BaAKCgLqvI02yZclWUnEWv+dEMw/P6Z6jWrHxEoT091VeXJpsWMgUjNZbG8AstiwewWwMFivMD83S9zBVc52Myol0q9ylbu1CH6rVkx6Cn0DRO1opI8niF3ay77/4mSJt9D8=</diagram></mxfile>">	1	<svg host="65bd71144e" xmlns="http://www.w3.org/2000/svg" style="background: transparent; background-color: transparent;" xmlns:xlink="http://www.w3.org/1999/xlink" version="1.1" width="1111px" height="601px" viewBox="-0.5 -0.5 1111 601" content="<mxfile><diagram id="uiDLYcUB2_ILZagOPzox" name="Page-1">7Vxbc5s4FP41THd3Jh7ul0d8SZuZps0k6Wz75MEg29oC8mI5sffXrwTCIBAOtnGcZMxDgo4uCH0fR+ccSZa0QbT+nHiL+S0KQCipcrCWtKGkqoqlkL9UsMkEhmZlglkCA1amEDzA/wATyky6ggFYcgUxQiGGC17oozgGPuZkXpKgZ77YFIX8UxfeDNQED74X1qV/wwDPM6ltyIX8C4Czef5kRWY5kZcXZoLl3AvQc0mkjSRtkCCEs7toPQAhHbt8XLJ61w25244lIMZtKqhZhScvXLF3Y/3Cm/xlE7SKA0DLy5LWf55DDB4Wnk9znwm6RDbHUUhSCrn1QjiLyb1Png8SIngCCYZk5FyWMUEYo4hkTFGMr70IhpQDX0D4BGg5lsEgV1SWHqAQJWl3tOyij1phtMzK0Y4lgCS8Sdpr2pMpDMNStev0IvIlTtBvUMox04vkhN4EhHekTQyR8B2+VgpgRN++PuQMBVoNrEsiBsFngCKAkw0pkufKOUE2PNGfC3YpSi6cl6hl5hU9RunZtvECdXLDgBeTQKthDgLCd5ZECZ6jGYq9cFRI+zwrijJfER2UFIF/AMYbhiQFi2cKGa1k85PW7xl58hdrLk0M11xqw1IH8SYAU28VYp41ioA1jWAu0Srx2egYTOV4yQzkpexMRgduJ+IJCD0Mn3hFcgx4+gW8o8FzzgWeIVC/ZvqqdAA4WM1/VyjPuMoGwSUFFH2xTl89zyd3s/T/yJJcU+qTG1NyDMnVpJEtuSOpfy2NSHIoOTYt01fSMhbNcqw8S877Qd4g6wpr9bSzQwSDIGVoa6LsNwsEHrCnft4MewtFOCv4NphMO1Lvms2rd1UW6HdVoN71DrS7dVEQeymIfJotawjVPJeGsC/oHY+e0TV6aVU3SbxNqcACwRgvSy3fUUFJDZgOrwZ0Xe8ZvHlerWJrxu4q5CbrR8Gm7Qu1IphzIdjxBJPfLcEM/dQEywdMYOQsF158lJETIi8gz19ilABSbvjwtWS3ZK1/CLtlavvAb2W3TGyCqNyR3WJqPJ80MyfHq1guilZjDg57DHNyx2CvIEteDh8KXxUG9sWRdo2+ZAzrKoQf/xjFoAJeLjpOcfGwd4CsUbFINU2ErCMAVusC2Ga/J4BPxTebiSZCFTHx/N+z9Cu+8rPxopoCxhBDLxQqCxwtym7NZPscbcho9QeMFziR1P7mz+3Dk2p3SNVSJyHRInRQ16SW/Vepyy88sz1rQzDFb5ezeiNnO1FAeiUuZgkcJ13AUqMLltr7T1wUxqtsAqF0jBGbTAT+eXeOvkMd975OHX17JLmGwNFv8uY/4mR5Jie/Olmq9is6+XkI681z1ZKc/oWr5+aqYp6Rq2qHDoHcnsQPN7ekQTkNh5pS36E3riLZ9sdk1plcBr3CLEMVugzWqbilCLhVjW/EgUsXgQtLJ/CW8xRYhQeRR/y0wQewhvgnexK9T8MhPcsxWLqIiNDEppS4AwkkA0VZVoRMsuCKQ9RpObxC2tOk3SGWNFVtspEZL0Y8SqAbO+y0/QIj9TBGdQHVkJWeyreSBXFYxR0hEdKY1bMsZ3tZXMu64/Sc8qXzj8niQrXHHBA4yalb4vIARQsYElhG11J/QHTXI1F6phdR5ZP9/Xb34+rmPs+vB/fer3fc7Gl07h3LFVPO0PW6CrNP5Byr9ajHTRysfEze+oLmIWgaFWPH0gRTkmhG6gROvQanS8YVr+KPFLp6RTjNiqrXFZGFcbKvszl0xRuPnTlfMA5gAnw8ZoHPbZqFPxuM1zfo/OUzVTxZLniPrpWjt1MsiBu+pXcnzscjcz5cm7od1AsZpm7vDndkv2HhhuDiu+yvWSzjvL6L+YJqOY1fnG4CciX3OuPoxRfubvlMb8Un+VR8OiB+vZtPDEiaw1YlhNOVYEFk3+np+H5s2lL5/Rpd7dYL5S6ovFV6+UKMIwgYiojcic1VD25f1oG7MaaV6qKFUcf1VKvAWj0O7GfxjXGEAl4fLGGExygON9yHXrX8aCnSYFZOxjCE8UxKj4FM4aylpfReF2JfkTSmWvHA9Fe1k7R6jHc3bzCIFqGHQRvuFGUvtOmYNrbeijYiddMJbepxtV20CV5iS3BRMafjil4N8oi5Itr60QlX6kG7IgYrtzpa8XHMEfMVYa9sYBVYI4poGakTc8SqYe7S2gRd8sf4QIC+5j5DpQKoVV9JOVXkXav7vxcYD7T4TKMKo0Adi7RxJ0DW/b8LkAfaYMo5gdTrKrYGY/n0CRvE8sGRfJdEL98X8auU07BH4sDDJlnXKjscXj4fwrQOd8BUFqPS1XGQIzYZ6HVn6ubb8Mfg8fv9mNzc3I8Gj+Pb0e33+1/j2+9D0sbwEzV+x6lnFcH1p8uHeNBaQ8XAFW/A1072Jdb15+4PrxX3dbPO/dOcvavvANIrkSylugW8Yf/PId9M80KNYBmSrR3620WGwq3U5PRqXndsHR9Xdx2QKk55p7uCHNfDIO6BaAKCgLqvI02yZclWUnEWv+dEMw/P6Z6jWrHxEoT091VeXJpsWMgUjNZbG8AstiwewWwMFivMD83S9zBVc52Myol0q9ylbu1CH6rVkx6Cn0DRO1opI8niF3ay77/4mSJt9D8=</diagram></mxfile>">
2	<defs/>	2	<defs/>
3	<g>	3	<g>
4	<g>	4	<g>
5	<rect x="0" y="0" width="1110" height="600" fill="#ffffff" stroke="#666666" pointer-events="all" style="fill: light-dark(rgb(255, 255, 255), rgb(18, 18, 18)); stroke: light-dark(rgb(102, 102, 102), rgb(149, 149, 149));"/>	5	<rect x="0" y="0" width="1110" height="600" fill="#ffffff" stroke="#666666" pointer-events="all" style="fill: light-dark(rgb(255, 255, 255), rgb(18, 18, 18)); stroke: light-dark(rgb(102, 102, 102), rgb(149, 149, 149));"/>
6	</g>	6	</g>
7	<g>	7	<g>
8	<path d="M 500 210 L 690 210 L 690 263.63" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	8	<path d="M 500 210 L 690 210 L 690 263.63" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
9	<path d="M 690 268.88 L 686.5 261.88 L 690 263.63 L 693.5 261.88 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	9	<path d="M 690 268.88 L 686.5 261.88 L 690 263.63 L 693.5 261.88 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
10	</g>	10	</g>
11	<g>	11	<g>
12	<path d="M 380 210 L 220 210 L 220 263.63" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	12	<path d="M 380 210 L 220 210 L 220 263.63" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
13	<path d="M 220 268.88 L 216.5 261.88 L 220 263.63 L 223.5 261.88 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	13	<path d="M 220 268.88 L 216.5 261.88 L 220 263.63 L 223.5 261.88 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
14	</g>	14	</g>
15	<g>	15	<g>
16	<rect x="380" y="190" width="120" height="40" fill="#dae8fc" stroke="#6c8ebf" pointer-events="all" style="fill: light-dark(rgb(218, 232, 252), rgb(29, 41, 59)); stroke: light-dark(rgb(108, 142, 191), rgb(92, 121, 163));"/>	16	<rect x="380" y="190" width="120" height="40" fill="#dae8fc" stroke="#6c8ebf" pointer-events="all" style="fill: light-dark(rgb(218, 232, 252), rgb(29, 41, 59)); stroke: light-dark(rgb(108, 142, 191), rgb(92, 121, 163));"/>
17	</g>	17	</g>
18	<g>	18	<g>
19	<g transform="translate(-0.5 -0.5)">	19	<g transform="translate(-0.5 -0.5)">
20	<switch>	20	<switch>
21	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	21	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
22	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 210px; margin-left: 381px;">	22	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 210px; margin-left: 381px;">
23	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	23	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
24	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	24	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
25	<font style="font-size: 14px;">	25	<font style="font-size: 14px;">
26	离散访存类算子	26	离散访存类算子
27	</font>	27	</font>
28	</div>	28	</div>
29	</div>	29	</div>
30	</div>	30	</div>
31	</foreignObject>	31	</foreignObject>
32	<text x="440" y="214" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	32	<text x="440" y="214" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
33	离散访存类算子	33	离散访存类算子
34	</text>	34	</text>
35	</switch>	35	</switch>
36	</g>	36	</g>
37	</g>	37	</g>
38	<g>	38	<g>
39	<path d="M 690 394.5 L 690 508.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	39	<path d="M 690 394.5 L 690 508.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
40	<path d="M 690 513.38 L 686.5 506.38 L 690 508.13 L 693.5 506.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	40	<path d="M 690 513.38 L 686.5 506.38 L 690 508.13 L 693.5 506.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
41	</g>	41	</g>
42	<g>	42	<g>
43	<path d="M 690 394.5 L 690 434.5 L 835 434.5 L 835 508.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	43	<path d="M 690 394.5 L 690 434.5 L 835 434.5 L 835 508.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
44	<path d="M 835 513.38 L 831.5 506.38 L 835 508.13 L 838.5 506.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	44	<path d="M 835 513.38 L 831.5 506.38 L 835 508.13 L 838.5 506.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
45	</g>	45	</g>
46	<g>	46	<g>
47	<path d="M 690 394.5 L 690 434.5 L 545 434.5 L 545 508.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	47	<path d="M 690 394.5 L 690 434.5 L 545 434.5 L 545 508.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
48	<path d="M 545 513.38 L 541.5 506.38 L 545 508.13 L 548.5 506.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	48	<path d="M 545 513.38 L 541.5 506.38 L 545 508.13 L 548.5 506.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
49	</g>	49	</g>
50	<g>	50	<g>
51	<rect x="630" y="354.5" width="120" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>	51	<rect x="630" y="354.5" width="120" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
52	</g>	52	</g>
53	<g>	53	<g>
54	<g transform="translate(-0.5 -0.5)">	54	<g transform="translate(-0.5 -0.5)">
55	<switch>	55	<switch>
56	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	56	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
57	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 375px; margin-left: 631px;">	57	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 375px; margin-left: 631px;">
58	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	58	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
59	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	59	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
60	<span style="font-size: 14px;">	60	<span style="font-size: 14px;">
61	load/store DSL	61	load/store DSL
62	</span>	62	</span>
63	</div>	63	</div>
64	</div>	64	</div>
65	</div>	65	</div>
66	</foreignObject>	66	</foreignObject>
67	<text x="690" y="378" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	67	<text x="690" y="378" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
68	load/store DSL	68	load/store DSL
69	</text>	69	</text>
70	</switch>	70	</switch>
71	</g>	71	</g>
72	</g>	72	</g>
73	<g>	73	<g>
74	<rect x="580" y="324.5" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>	74	<rect x="580" y="324.5" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>
75	</g>	75	</g>
76	<g>	76	<g>
77	<g transform="translate(-0.5 -0.5)">	77	<g transform="translate(-0.5 -0.5)">
78	<switch>	78	<switch>
79	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	79	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
80	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 340px; margin-left: 625px;">	80	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 340px; margin-left: 625px;">
81	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	81	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
82	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	82	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
83	tl.load/tl.store	83	tl.load/tl.store
84	</div>	84	</div>
85	</div>	85	</div>
86	</div>	86	</div>
87	</foreignObject>	87	</foreignObject>
88	<text x="625" y="343" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	88	<text x="625" y="343" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
89	tl.load/tl.store	89	tl.load/tl.store
90	</text>	90	</text>
91	</switch>	91	</switch>
92	</g>	92	</g>
93	</g>	93	</g>
94	<g>	94	<g>
95	<rect x="640" y="160" width="140" height="50" fill="none" stroke="none" pointer-events="all"/>	95	<rect x="640" y="160" width="140" height="50" fill="none" stroke="none" pointer-events="all"/>
96	</g>	96	</g>
97	<g>	97	<g>
98	<g transform="translate(-0.5 -0.5)">	98	<g transform="translate(-0.5 -0.5)">
99	<switch>	99	<switch>
100	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	100	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
101	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 185px; margin-left: 642px;">	101	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 185px; margin-left: 642px;">
102	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">	102	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">
103	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">	103	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
104	<div>	104	<div>
105	<b style="background-color: initial;">	105	<b style="background-color: initial;">
106	tmp0	106	tmp0
107	</b>	107	</b>
108	=tl.load(inptr+y)	108	=tl.load(inptr+y)
109	<br/>	109	<br/>
110	</div>	110	</div>
111	index=x+8*	111	index=x+8*
112	<b>	112	<b>
113	tmp0	113	tmp0
114	</b>	114	</b>
115	</div>	115	</div>
116	</div>	116	</div>
117	</div>	117	</div>
118	</foreignObject>	118	</foreignObject>
119	<text x="642" y="189" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px">	119	<text x="642" y="189" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px">
120	tmp0=tl.load(inptr+y...	120	tmp0=tl.load(inptr+y...
121	</text>	121	</text>
122	</switch>	122	</switch>
123	</g>	123	</g>
124	</g>	124	</g>
125	<g>	125	<g>
126	<rect x="630" y="270" width="120" height="40" fill="#dae8fc" stroke="#6c8ebf" pointer-events="all" style="fill: light-dark(rgb(218, 232, 252), rgb(29, 41, 59)); stroke: light-dark(rgb(108, 142, 191), rgb(92, 121, 163));"/>	126	<rect x="630" y="270" width="120" height="40" fill="#dae8fc" stroke="#6c8ebf" pointer-events="all" style="fill: light-dark(rgb(218, 232, 252), rgb(29, 41, 59)); stroke: light-dark(rgb(108, 142, 191), rgb(92, 121, 163));"/>
127	</g>	127	</g>
128	<g>	128	<g>
129	<g transform="translate(-0.5 -0.5)">	129	<g transform="translate(-0.5 -0.5)">
130	<switch>	130	<switch>
131	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	131	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
132	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 290px; margin-left: 631px;">	132	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 290px; margin-left: 631px;">
133	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	133	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
134	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	134	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
135	<span style="text-wrap: nowrap;">	135	<span style="text-wrap: nowrap;">
136	<font style="font-size: 14px;">	136	<font style="font-size: 14px;">
137	间接访存	137	间接访存
138	</font>	138	</font>
139	</span>	139	</span>
140	</div>	140	</div>
141	</div>	141	</div>
142	</div>	142	</div>
143	</foreignObject>	143	</foreignObject>
144	<text x="690" y="294" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	144	<text x="690" y="294" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
145	间接访存	145	间接访存
146	</text>	146	</text>
147	</switch>	147	</switch>
148	</g>	148	</g>
149	</g>	149	</g>
150	<g>	150	<g>
151	<rect x="160" y="270" width="120" height="40" fill="#dae8fc" stroke="#6c8ebf" pointer-events="all" style="fill: light-dark(rgb(218, 232, 252), rgb(29, 41, 59)); stroke: light-dark(rgb(108, 142, 191), rgb(92, 121, 163));"/>	151	<rect x="160" y="270" width="120" height="40" fill="#dae8fc" stroke="#6c8ebf" pointer-events="all" style="fill: light-dark(rgb(218, 232, 252), rgb(29, 41, 59)); stroke: light-dark(rgb(108, 142, 191), rgb(92, 121, 163));"/>
152	</g>	152	</g>
153	<g>	153	<g>
154	<g transform="translate(-0.5 -0.5)">	154	<g transform="translate(-0.5 -0.5)">
155	<switch>	155	<switch>
156	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	156	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
157	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 290px; margin-left: 161px;">	157	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 118px; height: 1px; padding-top: 290px; margin-left: 161px;">
158	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	158	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
159	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	159	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
160	<span style="text-wrap: nowrap;">	160	<span style="text-wrap: nowrap;">
161	<font style="font-size: 14px;">	161	<font style="font-size: 14px;">
162	直接访存	162	直接访存
163	</font>	163	</font>
164	</span>	164	</span>
165	</div>	165	</div>
166	</div>	166	</div>
167	</div>	167	</div>
168	</foreignObject>	168	</foreignObject>
169	<text x="220" y="294" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	169	<text x="220" y="294" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
170	直接访存	170	直接访存
171	</text>	171	</text>
172	</switch>	172	</switch>
173	</g>	173	</g>
174	</g>	174	</g>
175	<g>	175	<g>
176	<rect x="460" y="514.5" width="170" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>	176	<rect x="460" y="514.5" width="170" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
177	</g>	177	</g>
178	<g>	178	<g>
179	<g transform="translate(-0.5 -0.5)">	179	<g transform="translate(-0.5 -0.5)">
180	<switch>	180	<switch>
181	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	181	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
182	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 535px; margin-left: 461px;">	182	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 535px; margin-left: 461px;">
183	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	183	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
184	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	184	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
185	<span style="font-size: 14px; text-wrap: nowrap;">	185	<span style="font-size: 14px; text-wrap: nowrap;">
186	SIMD 方案	186	SIMD 方案
187	</span>	187	</span>
188	</div>	188	</div>
189	</div>	189	</div>
190	</div>	190	</div>
191	</foreignObject>	191	</foreignObject>
192	<text x="545" y="538" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	192	<text x="545" y="538" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
193	SIMD 方案	193	SIMD 方案
194	</text>	194	</text>
195	</switch>	195	</switch>
196	</g>	196	</g>
197	</g>	197	</g>
198	<g>	198	<g>
199	<path d="M 0 491.2 L 1107.78 490" fill="none" stroke="#000000" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	199	<path d="M 0 491.2 L 1107.78 490" fill="none" stroke="#000000" stroke-miterlimit="10" stroke-dasharray="3 3" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
200	</g>	200	</g>
201	<g>	201	<g>
202	<rect x="30" y="534" width="180" height="30" fill="none" stroke="none" pointer-events="all"/>	202	<rect x="30" y="534" width="180" height="30" fill="none" stroke="none" pointer-events="all"/>
203	</g>	203	</g>
204	<g>	204	<g>
205	<g transform="translate(-0.5 -0.5)">	205	<g transform="translate(-0.5 -0.5)">
206	<switch>	206	<switch>
207	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	207	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
208	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 549px; margin-left: 120px;">	208	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 549px; margin-left: 120px;">
209	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	209	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
210	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	210	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
211	Compiler（TA&NPU-IR）	211	Compiler（TA&NPU-IR）
212	</div>	212	</div>
213	</div>	213	</div>
214	</div>	214	</div>
215	</foreignObject>	215	</foreignObject>
216	<text x="120" y="553" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px" text-anchor="middle" font-weight="bold">	216	<text x="120" y="553" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px" text-anchor="middle" font-weight="bold">
217	Compiler（TA&NPU-IR）	217	Compiler（TA&NPU-IR）
218	</text>	218	</text>
219	</switch>	219	</switch>
220	</g>	220	</g>
221	</g>	221	</g>
222	<g>	222	<g>
223	<rect x="50" y="263.5" width="70" height="30" fill="none" stroke="none" pointer-events="all"/>	223	<rect x="50" y="263.5" width="70" height="30" fill="none" stroke="none" pointer-events="all"/>
224	</g>	224	</g>
225	<g>	225	<g>
226	<g transform="translate(-0.5 -0.5)">	226	<g transform="translate(-0.5 -0.5)">
227	<switch>	227	<switch>
228	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	228	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
229	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 279px; margin-left: 85px;">	229	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 279px; margin-left: 85px;">
230	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	230	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
231	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	231	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
232	Inductor	232	Inductor
233	</div>	233	</div>
234	</div>	234	</div>
235	</div>	235	</div>
236	</foreignObject>	236	</foreignObject>
237	<text x="85" y="283" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px" text-anchor="middle" font-weight="bold">	237	<text x="85" y="283" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px" text-anchor="middle" font-weight="bold">
238	Inductor	238	Inductor
239	</text>	239	</text>
240	</switch>	240	</switch>
241	</g>	241	</g>
242	</g>	242	</g>
243	<g>	243	<g>
244	<rect x="600" y="404.5" width="80" height="30" fill="none" stroke="none" pointer-events="all"/>	244	<rect x="600" y="404.5" width="80" height="30" fill="none" stroke="none" pointer-events="all"/>
245	</g>	245	</g>
246	<g>	246	<g>
247	<g transform="translate(-0.5 -0.5)">	247	<g transform="translate(-0.5 -0.5)">
248	<switch>	248	<switch>
249	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	249	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
250	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 420px; margin-left: 640px;">	250	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 420px; margin-left: 640px;">
251	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	251	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
252	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	252	<div style="display: inline-block; font-size: 14px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
253	Autotune	253	Autotune
254	</div>	254	</div>
255	</div>	255	</div>
256	</div>	256	</div>
257	</foreignObject>	257	</foreignObject>
258	<text x="640" y="424" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px" text-anchor="middle" font-weight="bold">	258	<text x="640" y="424" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="14px" text-anchor="middle" font-weight="bold">
259	Autotune	259	Autotune
260	</text>	260	</text>
261	</switch>	261	</switch>
262	</g>	262	</g>
263	</g>	263	</g>
264	<g>	264	<g>
265	<rect x="750" y="514.5" width="170" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>	265	<rect x="750" y="514.5" width="170" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
266	</g>	266	</g>
267	<g>	267	<g>
268	<g transform="translate(-0.5 -0.5)">	268	<g transform="translate(-0.5 -0.5)">
269	<switch>	269	<switch>
270	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	270	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
271	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 535px; margin-left: 751px;">	271	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 168px; height: 1px; padding-top: 535px; margin-left: 751px;">
272	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	272	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
273	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	273	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
274	<span>	274	<span>
275	<span style="text-wrap: nowrap;">	275	<span style="text-wrap: nowrap;">
276	indirect_load/indirect_store	276	indirect_load/indirect_store
277	</span>	277	</span>
278	<span style="text-wrap: nowrap;">	278	<span style="text-wrap: nowrap;">
279	<font style="font-size: 14px;">	279	<font style="font-size: 14px;">
280	</font>	280	</font>
281	</span>	281	</span>
282	</span>	282	</span>
283	<div>	283	<div>
284	<span style="text-wrap: nowrap;">	284	<span style="text-wrap: nowrap;">
285	<font style="font-size: 14px;">	285	<font style="font-size: 14px;">
286	SIMT 模板方案	286	SIMT 模板方案
287	</font>	287	</font>
288	</span>	288	</span>
289	</div>	289	</div>
290	</div>	290	</div>
291	</div>	291	</div>
292	</div>	292	</div>
293	</foreignObject>	293	</foreignObject>
294	<text x="835" y="538" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	294	<text x="835" y="538" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
295	indirect_load/indirect_store...	295	indirect_load/indirect_store...
296	</text>	296	</text>
297	</switch>	297	</switch>
298	</g>	298	</g>
299	</g>	299	</g>
300	<g>	300	<g>
301	<rect x="640" y="514.5" width="100" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>	301	<rect x="640" y="514.5" width="100" height="40" fill="#f8cecc" stroke="#b85450" pointer-events="all" style="fill: light-dark(rgb(248, 206, 204), rgb(81, 45, 43)); stroke: light-dark(rgb(184, 84, 80), rgb(215, 129, 126));"/>
302	</g>	302	</g>
303	<g>	303	<g>
304	<g transform="translate(-0.5 -0.5)">	304	<g transform="translate(-0.5 -0.5)">
305	<switch>	305	<switch>
306	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	306	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
307	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 98px; height: 1px; padding-top: 535px; margin-left: 641px;">	307	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 98px; height: 1px; padding-top: 535px; margin-left: 641px;">
308	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	308	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
309	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">	309	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: normal; word-wrap: normal; ">
310	<span style="font-size: 14px; text-wrap: nowrap;">	310	<span style="font-size: 14px; text-wrap: nowrap;">
311	纯SIMT	311	纯SIMT
312	</span>	312	</span>
313	</div>	313	</div>
314	</div>	314	</div>
315	</div>	315	</div>
316	</foreignObject>	316	</foreignObject>
317	<text x="690" y="538" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	317	<text x="690" y="538" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
318	纯SIMT	318	纯SIMT
319	</text>	319	</text>
320	</switch>	320	</switch>
321	</g>	321	</g>
322	</g>	322	</g>
323	<g>	323	<g>
324	<rect x="170" y="180" width="100" height="30" fill="none" stroke="none" pointer-events="all"/>	324	<rect x="170" y="180" width="100" height="30" fill="none" stroke="none" pointer-events="all"/>
325	</g>	325	</g>
326	<g>	326	<g>
327	<g transform="translate(-0.5 -0.5)">	327	<g transform="translate(-0.5 -0.5)">
328	<switch>	328	<switch>
329	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	329	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
330	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 195px; margin-left: 220px;">	330	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 195px; margin-left: 220px;">
331	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	331	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
332	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">	332	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
333	<span style="font-size: 14px; text-align: left;">	333	<span style="font-size: 14px; text-align: left;">
334	index=x+8*	334	index=x+8*
335	</span>	335	</span>
336	<span style="font-size: 14px; text-align: left;">	336	<span style="font-size: 14px; text-align: left;">
337	y	337	y
338	</span>	338	</span>
339	</div>	339	</div>
340	</div>	340	</div>
341	</div>	341	</div>
342	</foreignObject>	342	</foreignObject>
343	<text x="220" y="199" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle">	343	<text x="220" y="199" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle">
344	index=x+8*y	344	index=x+8*y
345	</text>	345	</text>
346	</switch>	346	</switch>
347	</g>	347	</g>
348	</g>	348	</g>
349	<g>	349	<g>
350	<rect x="130" y="240" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>	350	<rect x="130" y="240" width="90" height="30" fill="none" stroke="none" pointer-events="all"/>
351	</g>	351	</g>
352	<g>	352	<g>
353	<g transform="translate(-0.5 -0.5)">	353	<g transform="translate(-0.5 -0.5)">
354	<switch>	354	<switch>
355	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	355	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
356	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 255px; margin-left: 175px;">	356	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 255px; margin-left: 175px;">
357	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	357	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
358	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	358	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
359	tl.load/tl.store	359	tl.load/tl.store
360	</div>	360	</div>
361	</div>	361	</div>
362	</div>	362	</div>
363	</foreignObject>	363	</foreignObject>
364	<text x="175" y="259" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	364	<text x="175" y="259" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
365	tl.load/tl.store	365	tl.load/tl.store
366	</text>	366	</text>
367	</switch>	367	</switch>
368	</g>	368	</g>
369	</g>	369	</g>
370	<g>	370	<g>
371	<rect x="620" y="434.5" width="170" height="40" fill="none" stroke="none" pointer-events="all"/>	371	<rect x="620" y="434.5" width="170" height="40" fill="none" stroke="none" pointer-events="all"/>
372	</g>	372	</g>
373	<g>	373	<g>
374	<g transform="translate(-0.5 -0.5)">	374	<g transform="translate(-0.5 -0.5)">
375	<switch>	375	<switch>
376	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	376	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
377	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 455px; margin-left: 622px;">	377	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 455px; margin-left: 622px;">
378	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">	378	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">
379	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	379	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
380	compile_mode="simt_only"	380	compile_mode="simt_only"
381	<div>	381	<div>
382	simt only tiling config	382	simt only tiling config
383	</div>	383	</div>
384	</div>	384	</div>
385	</div>	385	</div>
386	</div>	386	</div>
387	</foreignObject>	387	</foreignObject>
388	<text x="622" y="458" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" font-weight="bold">	388	<text x="622" y="458" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" font-weight="bold">
389	compile_mode="simt_only"...	389	compile_mode="simt_only"...
390	</text>	390	</text>
391	</switch>	391	</switch>
392	</g>	392	</g>
393	</g>	393	</g>
394	<g>	394	<g>
395	<rect x="840" y="434.5" width="190" height="40" fill="none" stroke="none" pointer-events="all"/>	395	<rect x="840" y="434.5" width="190" height="40" fill="none" stroke="none" pointer-events="all"/>
396	</g>	396	</g>
397	<g>	397	<g>
398	<g transform="translate(-0.5 -0.5)">	398	<g transform="translate(-0.5 -0.5)">
399	<switch>	399	<switch>
400	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	400	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
401	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 455px; margin-left: 842px;">	401	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 455px; margin-left: 842px;">
402	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">	402	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">
403	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	403	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
404	compile_mode="simt_template"	404	compile_mode="simt_template"
405	<div>	405	<div>
406	simt template config	406	simt template config
407	</div>	407	</div>
408	</div>	408	</div>
409	</div>	409	</div>
410	</div>	410	</div>
411	</foreignObject>	411	</foreignObject>
412	<text x="842" y="458" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" font-weight="bold">	412	<text x="842" y="458" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" font-weight="bold">
413	compile_mode="simt_template"...	413	compile_mode="simt_template"...
414	</text>	414	</text>
415	</switch>	415	</switch>
416	</g>	416	</g>
417	</g>	417	</g>
418	<g>	418	<g>
419	<rect x="400" y="434.5" width="140" height="40" fill="none" stroke="none" pointer-events="all"/>	419	<rect x="400" y="434.5" width="140" height="40" fill="none" stroke="none" pointer-events="all"/>
420	</g>	420	</g>
421	<g>	421	<g>
422	<g transform="translate(-0.5 -0.5)">	422	<g transform="translate(-0.5 -0.5)">
423	<switch>	423	<switch>
424	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	424	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
425	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 455px; margin-left: 402px;">	425	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe flex-start; width: 1px; height: 1px; padding-top: 455px; margin-left: 402px;">
426	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">	426	<div style="box-sizing: border-box; font-size: 0; text-align: left; color: #000000; ">
427	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	427	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
428	compile_mode="simd"	428	compile_mode="simd"
429	<div>	429	<div>
430	simd tiling config	430	simd tiling config
431	</div>	431	</div>
432	</div>	432	</div>
433	</div>	433	</div>
434	</div>	434	</div>
435	</foreignObject>	435	</foreignObject>
436	<text x="402" y="458" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" font-weight="bold">	436	<text x="402" y="458" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" font-weight="bold">
437	compile_mode="simd"...	437	compile_mode="simd"...
438	</text>	438	</text>
439	</switch>	439	</switch>
440	</g>	440	</g>
441	</g>	441	</g>
442	<g>	442	<g>
443	<rect x="490" y="40" width="150" height="30" fill="none" stroke="none" pointer-events="all"/>	443	<rect x="490" y="40" width="150" height="30" fill="none" stroke="none" pointer-events="all"/>
444	</g>	444	</g>
445	<g>	445	<g>
446	<g transform="translate(-0.5 -0.5)">	446	<g transform="translate(-0.5 -0.5)">
447	<switch>	447	<switch>
448	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	448	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
449	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 55px; margin-left: 565px;">	449	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 55px; margin-left: 565px;">
450	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	450	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
451	<div style="display: inline-block; font-size: 16px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	451	<div style="display: inline-block; font-size: 16px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
452	Inductor 离散访存	452	Inductor 离散访存
453	</div>	453	</div>
454	</div>	454	</div>
455	</div>	455	</div>
456	</foreignObject>	456	</foreignObject>
457	<text x="565" y="60" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="16px" text-anchor="middle" font-weight="bold">	457	<text x="565" y="60" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="16px" text-anchor="middle" font-weight="bold">
458	Inductor 离散访存	458	Inductor 离散访存
459	</text>	459	</text>
460	</switch>	460	</switch>
461	</g>	461	</g>
462	</g>	462	</g>
463	<g>	463	<g>
464	<rect x="510" y="564" width="70" height="30" fill="none" stroke="none" pointer-events="all"/>	464	<rect x="510" y="564" width="70" height="30" fill="none" stroke="none" pointer-events="all"/>
465	</g>	465	</g>
466	<g>	466	<g>
467	<g transform="translate(-0.5 -0.5)">	467	<g transform="translate(-0.5 -0.5)">
468	<switch>	468	<switch>
469	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	469	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
470	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 579px; margin-left: 545px;">	470	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 579px; margin-left: 545px;">
471	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	471	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
472	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	472	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
473	A2/A3/A5	473	A2/A3/A5
474	</div>	474	</div>
475	</div>	475	</div>
476	</div>	476	</div>
477	</foreignObject>	477	</foreignObject>
478	<text x="545" y="583" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	478	<text x="545" y="583" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
479	A2/A3/A5	479	A2/A3/A5
480	</text>	480	</text>
481	</switch>	481	</switch>
482	</g>	482	</g>
483	</g>	483	</g>
484	<g>	484	<g>
485	<rect x="665" y="564.5" width="40" height="30" fill="none" stroke="none" pointer-events="all"/>	485	<rect x="665" y="564.5" width="40" height="30" fill="none" stroke="none" pointer-events="all"/>
486	</g>	486	</g>
487	<g>	487	<g>
488	<g transform="translate(-0.5 -0.5)">	488	<g transform="translate(-0.5 -0.5)">
489	<switch>	489	<switch>
490	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	490	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
491	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 580px; margin-left: 685px;">	491	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 580px; margin-left: 685px;">
492	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	492	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
493	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	493	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
494	A5	494	A5
495	</div>	495	</div>
496	</div>	496	</div>
497	</div>	497	</div>
498	</foreignObject>	498	</foreignObject>
499	<text x="685" y="583" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	499	<text x="685" y="583" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
500	A5	500	A5
501	</text>	501	</text>
502	</switch>	502	</switch>
503	</g>	503	</g>
504	</g>	504	</g>
505	<g>	505	<g>
506	<rect x="815" y="564.5" width="40" height="30" fill="none" stroke="none" pointer-events="all"/>	506	<rect x="815" y="564.5" width="40" height="30" fill="none" stroke="none" pointer-events="all"/>
507	</g>	507	</g>
508	<g>	508	<g>
509	<g transform="translate(-0.5 -0.5)">	509	<g transform="translate(-0.5 -0.5)">
510	<switch>	510	<switch>
511	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	511	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
512	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 580px; margin-left: 835px;">	512	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 580px; margin-left: 835px;">
513	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	513	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
514	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	514	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
515	A5	515	A5
516	</div>	516	</div>
517	</div>	517	</div>
518	</div>	518	</div>
519	</foreignObject>	519	</foreignObject>
520	<text x="835" y="583" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	520	<text x="835" y="583" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
521	A5	521	A5
522	</text>	522	</text>
523	</switch>	523	</switch>
524	</g>	524	</g>
525	</g>	525	</g>
526	<g>	526	<g>
527	<path d="M 690 310 L 690 348.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	527	<path d="M 690 310 L 690 348.13" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
528	<path d="M 690 353.38 L 686.5 346.38 L 690 348.13 L 693.5 346.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	528	<path d="M 690 353.38 L 686.5 346.38 L 690 348.13 L 693.5 346.38 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
529	</g>	529	</g>
530	<g>	530	<g>
531	<rect x="700" y="324.5" width="340" height="30" fill="none" stroke="none" pointer-events="all"/>	531	<rect x="700" y="324.5" width="340" height="30" fill="none" stroke="none" pointer-events="all"/>
532	</g>	532	</g>
533	<g>	533	<g>
534	<g transform="translate(-0.5 -0.5)">	534	<g transform="translate(-0.5 -0.5)">
535	<switch>	535	<switch>
536	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	536	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
537	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 340px; margin-left: 870px;">	537	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 340px; margin-left: 870px;">
538	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	538	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
539	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">	539	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; font-weight: bold; white-space: nowrap; ">
540	INDUCTOR_INDIRECT_MEMORY_MODE='simd_simt_mix'	540	INDUCTOR_INDIRECT_MEMORY_MODE='simd_simt_mix'
541	</div>	541	</div>
542	</div>	542	</div>
543	</div>	543	</div>
544	</foreignObject>	544	</foreignObject>
545	<text x="870" y="343" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">	545	<text x="870" y="343" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle" font-weight="bold">
546	INDUCTOR_INDIRECT_MEMORY_MODE='simd_simt_mix'	546	INDUCTOR_INDIRECT_MEMORY_MODE='simd_simt_mix'
547	</text>	547	</text>
548	</switch>	548	</switch>
549	</g>	549	</g>
550	</g>	550	</g>
551	<g>	551	<g>
552	<path d="M 440 140 L 440 183.63" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	552	<path d="M 440 140 L 440 183.63" fill="none" stroke="#000000" stroke-miterlimit="10" pointer-events="stroke" style="stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
553	<path d="M 440 188.88 L 436.5 181.88 L 440 183.63 L 443.5 181.88 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>	553	<path d="M 440 188.88 L 436.5 181.88 L 440 183.63 L 443.5 181.88 Z" fill="#000000" stroke="#000000" stroke-miterlimit="10" pointer-events="all" style="fill: light-dark(rgb(0, 0, 0), rgb(255, 255, 255)); stroke: light-dark(rgb(0, 0, 0), rgb(255, 255, 255));"/>
554	</g>	554	</g>
555	<g>	555	<g>
556	<rect x="240" y="100" width="400" height="40" fill="none" stroke="none" pointer-events="all"/>	556	<rect x="240" y="100" width="400" height="40" fill="none" stroke="none" pointer-events="all"/>
557	</g>	557	</g>
558	<g>	558	<g>
559	<g transform="translate(-0.5 -0.5)">	559	<g transform="translate(-0.5 -0.5)">
560	<switch>	560	<switch>
561	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">	561	<foreignObject style="overflow: visible; text-align: left;" pointer-events="none" width="100%" height="100%" requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility">
562	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 120px; margin-left: 440px;">	562	<div xmlns="http://www.w3.org/1999/xhtml" style="display: flex; align-items: unsafe center; justify-content: unsafe center; width: 1px; height: 1px; padding-top: 120px; margin-left: 440px;">
563	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">	563	<div style="box-sizing: border-box; font-size: 0; text-align: center; color: #000000; ">
564	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">	564	<div style="display: inline-block; font-size: 12px; font-family: "Helvetica"; color: light-dark(#000000, #ffffff); line-height: 1.2; pointer-events: all; white-space: nowrap; ">
565	<div>	565	<div>
566	<font color="#000000" style="color: light-dark(rgb(0, 0, 0), rgb(237, 237, 237));">	566	<font color="#000000" style="color: light-dark(rgb(0, 0, 0), rgb(237, 237, 237));">
567	<span style="font-size: 12px;">	567	<span style="font-size: 12px;">
568	load类：aten.embedding、aten.index、aten.gather、aten.index_select	568	load类：aten.embedding、aten.index、aten.gather、aten.index_select
569	</span>	569	</span>
570	</font>	570	</font>
571	</div>	571	</div>
572	<div>	572	<div>
573	<font color="#000000" style="color: light-dark(rgb(0, 0, 0), rgb(237, 237, 237));">	573	<font color="#000000" style="color: light-dark(rgb(0, 0, 0), rgb(237, 237, 237));">
574	<span style="font-size: 12px;">	574	<span style="font-size: 12px;">
575	store类：aten.index_put、aten.scatter	575	store类：aten.index_put、aten.scatter
576	</span>	576	</span>
577	</font>	577	</font>
578	</div>	578	</div>
579	</div>	579	</div>
580	</div>	580	</div>
581	</div>	581	</div>
582	</foreignObject>	582	</foreignObject>
583	<text x="440" y="124" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle">	583	<text x="440" y="124" fill="light-dark(#000000, #ffffff)" font-family=""Helvetica"" font-size="12px" text-anchor="middle">
584	load类：aten.embedding、aten.index、aten.gather、aten.index_select...	584	load类：aten.embedding、aten.index、aten.gather、aten.index_select...
585	</text>	585	</text>
586	</switch>	586	</switch>
587	</g>	587	</g>
588	</g>	588	</g>
589	</g>	589	</g>
590	<switch>	590	<switch>
591	<g requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility"/>	591	<g requiredFeatures="http://www.w3.org/TR/SVG11/feature#Extensibility"/>
592	<a transform="translate(0,-5)" xlink:href="https://www.drawio.com/doc/faq/svg-export-text-problems" target="_blank">	592	<a transform="translate(0,-5)" xlink:href="https://www.drawio.com/doc/faq/svg-export-text-problems" target="_blank">
593	<text text-anchor="middle" font-size="10px" x="50%" y="100%">	593	<text text-anchor="middle" font-size="10px" x="50%" y="100%">
594	Text is not SVG - cannot display	594	Text is not SVG - cannot display
595	</text>	595	</text>
596	</a>	596	</a>
597	</switch>	597	</switch>
598	</svg>	598	</svg>

Mtorch_npu/_inductor/docs/profiling/trace.drawio.svg+7-7

此文件变更行数或变更字符数较多，你可以直接查看源码

Mtorch_npu/_inductor/fx_passes/pattern_match/__init__.py+3-3

	@@ -1,3 +1,3 @@
1	from .npu_fusion_attention_graph import npu_fusion_attention_graph	1	from .npu_fusion_attention_graph import npu_fusion_attention_graph
2		2
3	__all__ = ["npu_fusion_attention_graph"]	3	__all__ = ["npu_fusion_attention_graph"]

Mtorch_npu/_inductor/fx_passes/pattern_match/npu_fusion_attention_graph.py+0-1

	@@ -254,4 +254,3 @@ def register_fa_pass():
254	register_replacement(	254	register_replacement(
255	**register_replacement_kwargs,	255	**register_replacement_kwargs,
256	)	256	)
257

Mtorch_npu/_inductor/lowering_fallback_list.py+1104-1104

	@@ -1,1105 +1,1105 @@
1	"""	1	"""
2	This file holds FALLBACK_LIST in which aten/primes ops will be fallbacked to AclNN op and	2	This file holds FALLBACK_LIST in which aten/primes ops will be fallbacked to AclNN op and
3	will not be allowed to lowering and fusion.	3	will not be allowed to lowering and fusion.
4		4
5	After fixed and verified, it can be removed from FALLBACK_LIST.	5	After fixed and verified, it can be removed from FALLBACK_LIST.
6	"""	6	"""
7		7
8	import torch	8	import torch
9		9
10	from torch.ops import (	10	from torch.ops import (
11	quantized_decomposed,	11	quantized_decomposed,
12	rngprims,	12	rngprims,
13	_inductor_test,	13	_inductor_test,
14	inductor,	14	inductor,
15	fsdp,	15	fsdp,
16	_c10d_functional,	16	_c10d_functional,
17	_dtensor,	17	_dtensor,
18	_quantized,	18	_quantized,
19	quantized,	19	quantized,
20	)	20	)
21		21
22	from torch._prims.rng_prims import (	22	from torch._prims.rng_prims import (
23	graphsafe_run_with_rng_state,	23	graphsafe_run_with_rng_state,
24	run_and_save_rng_state,	24	run_and_save_rng_state,
25	run_with_rng_state,	25	run_with_rng_state,
26	)	26	)
27	from torch._higher_order_ops.auto_functionalize import auto_functionalized	27	from torch._higher_order_ops.auto_functionalize import auto_functionalized
28	from torch._higher_order_ops.cond import cond	28	from torch._higher_order_ops.cond import cond
29	from torch._higher_order_ops.while_loop import while_loop	29	from torch._higher_order_ops.while_loop import while_loop
30	from torch._higher_order_ops._invoke_quant import invoke_quant	30	from torch._higher_order_ops._invoke_quant import invoke_quant
31	from torch._higher_order_ops.associative_scan import associative_scan	31	from torch._higher_order_ops.associative_scan import associative_scan
32	from torch._higher_order_ops.effects import with_effects	32	from torch._higher_order_ops.effects import with_effects
33		33
34	from .config import inductor_indirect_memory_mode, inductor_ascend_linear_mode	34	from .config import inductor_indirect_memory_mode, inductor_ascend_linear_mode
35		35
36	aten = torch.ops.aten	36	aten = torch.ops.aten
37	prims = torch.ops.prims	37	prims = torch.ops.prims
38		38
39		39
40	NPU_EXTRA_FALLBACK_LIST = [	40	NPU_EXTRA_FALLBACK_LIST = [
41	_c10d_functional.all_gather_into_tensor,	41	_c10d_functional.all_gather_into_tensor,
42	_c10d_functional.all_gather_into_tensor.default,	42	_c10d_functional.all_gather_into_tensor.default,
43	_c10d_functional.all_gather_into_tensor_coalesced,	43	_c10d_functional.all_gather_into_tensor_coalesced,
44	_c10d_functional.all_gather_into_tensor_coalesced.default,	44	_c10d_functional.all_gather_into_tensor_coalesced.default,
45	_c10d_functional.all_gather_into_tensor_out,	45	_c10d_functional.all_gather_into_tensor_out,
46	_c10d_functional.all_gather_into_tensor_out.default,	46	_c10d_functional.all_gather_into_tensor_out.default,
47	_c10d_functional.all_reduce,	47	_c10d_functional.all_reduce,
48	_c10d_functional.all_reduce.default,	48	_c10d_functional.all_reduce.default,
49	_c10d_functional.all_reduce_,	49	_c10d_functional.all_reduce_,
50	_c10d_functional.all_reduce_.default,	50	_c10d_functional.all_reduce_.default,
51	_c10d_functional.all_reduce_coalesced,	51	_c10d_functional.all_reduce_coalesced,
52	_c10d_functional.all_reduce_coalesced.default,	52	_c10d_functional.all_reduce_coalesced.default,
53	_c10d_functional.all_reduce_coalesced_,	53	_c10d_functional.all_reduce_coalesced_,
54	_c10d_functional.all_reduce_coalesced_.default,	54	_c10d_functional.all_reduce_coalesced_.default,
55	_c10d_functional.all_to_all_single,	55	_c10d_functional.all_to_all_single,
56	_c10d_functional.all_to_all_single.default,	56	_c10d_functional.all_to_all_single.default,
57	_c10d_functional.broadcast,	57	_c10d_functional.broadcast,
58	_c10d_functional.broadcast.default,	58	_c10d_functional.broadcast.default,
59	_c10d_functional.broadcast_,	59	_c10d_functional.broadcast_,
60	_c10d_functional.broadcast_.default,	60	_c10d_functional.broadcast_.default,
61	_c10d_functional.reduce_scatter_tensor,	61	_c10d_functional.reduce_scatter_tensor,
62	_c10d_functional.reduce_scatter_tensor.default,	62	_c10d_functional.reduce_scatter_tensor.default,
63	_c10d_functional.reduce_scatter_tensor_coalesced,	63	_c10d_functional.reduce_scatter_tensor_coalesced,
64	_c10d_functional.reduce_scatter_tensor_coalesced.default,	64	_c10d_functional.reduce_scatter_tensor_coalesced.default,
65	_c10d_functional.wait_tensor,	65	_c10d_functional.wait_tensor,
66	_c10d_functional.wait_tensor.default,	66	_c10d_functional.wait_tensor.default,
67	_dtensor.shard_dim_alltoall,	67	_dtensor.shard_dim_alltoall,
68	_dtensor.shard_dim_alltoall.default,	68	_dtensor.shard_dim_alltoall.default,
69	_inductor_test.realize,	69	_inductor_test.realize,
70	_inductor_test.realize.default,	70	_inductor_test.realize.default,
71	associative_scan,	71	associative_scan,
72	aten.__and__,	72	aten.__and__,
73	aten.__and__.bool,	73	aten.__and__.bool,
74	aten.__and__.int,	74	aten.__and__.int,
75	aten.__iand__,	75	aten.__iand__,
76	aten.__iand__.Scalar,	76	aten.__iand__.Scalar,
77	aten.__iand__.Tensor,	77	aten.__iand__.Tensor,
78	aten.__ilshift__,	78	aten.__ilshift__,
79	aten.__ilshift__.Scalar,	79	aten.__ilshift__.Scalar,
80	aten.__ilshift__.Tensor,	80	aten.__ilshift__.Tensor,
81	aten.__ior__.Scalar,	81	aten.__ior__.Scalar,
82	aten.__irshift__.Scalar,	82	aten.__irshift__.Scalar,
83	aten.__ixor__,	83	aten.__ixor__,
84	aten.__ixor__.Scalar,	84	aten.__ixor__.Scalar,
85	aten.__ixor__.Tensor,	85	aten.__ixor__.Tensor,
86	aten.__lshift__,	86	aten.__lshift__,
87	aten.__lshift__.Scalar,	87	aten.__lshift__.Scalar,
88	aten.__lshift__.Scalar_out,	88	aten.__lshift__.Scalar_out,
89	aten.__lshift__.Tensor,	89	aten.__lshift__.Tensor,
90	aten.__lshift__.Tensor_out,	90	aten.__lshift__.Tensor_out,
91	aten.__lshift__.int,	91	aten.__lshift__.int,
92	aten.__or__,	92	aten.__or__,
93	aten.__or__.bool,	93	aten.__or__.bool,
94	aten.__or__.int,	94	aten.__or__.int,
95	aten.__rshift__,	95	aten.__rshift__,
96	aten.__rshift__.Scalar,	96	aten.__rshift__.Scalar,
97	aten.__rshift__.Scalar_out,	97	aten.__rshift__.Scalar_out,
98	aten.__rshift__.Tensor,	98	aten.__rshift__.Tensor,
99	aten.__rshift__.Tensor_out,	99	aten.__rshift__.Tensor_out,
100	aten.__rshift__.int,	100	aten.__rshift__.int,
101	aten.__xor__,	101	aten.__xor__,
102	aten.__xor__.bool,	102	aten.__xor__.bool,
103	aten.__xor__.int,	103	aten.__xor__.int,
104	aten._adaptive_avg_pool2d,	104	aten._adaptive_avg_pool2d,
105	aten._adaptive_avg_pool2d.default,	105	aten._adaptive_avg_pool2d.default,
106	aten._adaptive_avg_pool2d.out,	106	aten._adaptive_avg_pool2d.out,
107	aten._assert_tensor_metadata,	107	aten._assert_tensor_metadata,
108	aten._assert_tensor_metadata.default,	108	aten._assert_tensor_metadata.default,
109	aten._convolution,	109	aten._convolution,
110	aten._convolution.default,	110	aten._convolution.default,
111	aten._convolution.out,	111	aten._convolution.out,
112	aten._foobar,	112	aten._foobar,
113	aten._foobar.default,	113	aten._foobar.default,
114	aten._foobar.out,	114	aten._foobar.out,
115	aten._int_mm,	115	aten._int_mm,
116	aten._int_mm.default,	116	aten._int_mm.default,
117	aten._int_mm.out,	117	aten._int_mm.out,
118	aten._jagged_to_padded_dense_forward.default,	118	aten._jagged_to_padded_dense_forward.default,
119	aten._local_scalar_dense,	119	aten._local_scalar_dense,
120	aten._local_scalar_dense.default,	120	aten._local_scalar_dense.default,
121	aten._neg_view,	121	aten._neg_view,
122	aten._neg_view.default,	122	aten._neg_view.default,
123	aten._padded_dense_to_jagged_forward,	123	aten._padded_dense_to_jagged_forward,
124	aten._padded_dense_to_jagged_forward.default,	124	aten._padded_dense_to_jagged_forward.default,
125	aten._scaled_mm.default,	125	aten._scaled_mm.default,
126	aten._sparse_semi_structured_mm,	126	aten._sparse_semi_structured_mm,
127	aten._sparse_semi_structured_mm.default,	127	aten._sparse_semi_structured_mm.default,
128	aten._unsafe_index_put.hacked_twin,	128	aten._unsafe_index_put.hacked_twin,
129	aten._unsafe_masked_index,	129	aten._unsafe_masked_index,
130	aten._unsafe_masked_index.default,	130	aten._unsafe_masked_index.default,
131	aten._unsafe_masked_index_put_accumulate,	131	aten._unsafe_masked_index_put_accumulate,
132	aten._unsafe_view,	132	aten._unsafe_view,
133	aten._weight_int4pack_mm_for_cpu,	133	aten._weight_int4pack_mm_for_cpu,
134	aten._weight_int4pack_mm_for_cpu.default,	134	aten._weight_int4pack_mm_for_cpu.default,
135	aten._weight_int8pack_mm,	135	aten._weight_int8pack_mm,
136	aten._weight_int8pack_mm.default,	136	aten._weight_int8pack_mm.default,
137	aten.acos,	137	aten.acos,
138	aten.acos.Scalar,	138	aten.acos.Scalar,
139	aten.acos.complex,	139	aten.acos.complex,
140	aten.acos.default,	140	aten.acos.default,
141	aten.acos.float,	141	aten.acos.float,
142	aten.acos.int,	142	aten.acos.int,
143	aten.acos.out,	143	aten.acos.out,
144	aten.acosh,	144	aten.acosh,
145	aten.acosh.Scalar,	145	aten.acosh.Scalar,
146	aten.acosh.complex,	146	aten.acosh.complex,
147	aten.acosh.default,	147	aten.acosh.default,
148	aten.acosh.float,	148	aten.acosh.float,
149	aten.acosh.int,	149	aten.acosh.int,
150	aten.acosh.out,	150	aten.acosh.out,
151	aten.adaptive_max_pool2d,	151	aten.adaptive_max_pool2d,
152	aten.add_,	152	aten.add_,
153	aten.add_.Scalar,	153	aten.add_.Scalar,
154	aten.add_.Tensor,	154	aten.add_.Tensor,
155	aten.add_.t,	155	aten.add_.t,
156	aten.alias,	156	aten.alias,
157	aten.alias.default,	157	aten.alias.default,
158	aten.any,	158	aten.any,
159	aten.any.all_out,	159	aten.any.all_out,
160	aten.any.bool,	160	aten.any.bool,
161	aten.any.default,	161	aten.any.default,
162	aten.any.dim,	162	aten.any.dim,
163	aten.any.dimname_out,	163	aten.any.dimname_out,
164	aten.any.dims,	164	aten.any.dims,
165	aten.any.dims_out,	165	aten.any.dims_out,
166	aten.any.float,	166	aten.any.float,
167	aten.any.int,	167	aten.any.int,
168	aten.any.out,	168	aten.any.out,
169	aten.any.str,	169	aten.any.str,
170	aten.as_strided,	170	aten.as_strided,
171	aten.as_strided.default,	171	aten.as_strided.default,
172	aten.as_strided_,	172	aten.as_strided_,
173	aten.as_strided_.default,	173	aten.as_strided_.default,
174	aten.as_strided_copy,	174	aten.as_strided_copy,
175	aten.as_strided_copy.default,	175	aten.as_strided_copy.default,
176	aten.as_strided_copy.out,	176	aten.as_strided_copy.out,
177	aten.as_strided_scatter,	177	aten.as_strided_scatter,
178	aten.as_strided_scatter.default,	178	aten.as_strided_scatter.default,
179	aten.as_strided_scatter.out,	179	aten.as_strided_scatter.out,
180	aten.asin,	180	aten.asin,
181	aten.asin.Scalar,	181	aten.asin.Scalar,
182	aten.asin.complex,	182	aten.asin.complex,
183	aten.asin.default,	183	aten.asin.default,
184	aten.asin.float,	184	aten.asin.float,
185	aten.asin.int,	185	aten.asin.int,
186	aten.asin.out,	186	aten.asin.out,
187	aten.asinh,	187	aten.asinh,
188	aten.asinh.Scalar,	188	aten.asinh.Scalar,
189	aten.asinh.complex,	189	aten.asinh.complex,
190	aten.asinh.default,	190	aten.asinh.default,
191	aten.asinh.float,	191	aten.asinh.float,
192	aten.asinh.int,	192	aten.asinh.int,
193	aten.asinh.out,	193	aten.asinh.out,
194	aten.atan,	194	aten.atan,
195	aten.atan.Scalar,	195	aten.atan.Scalar,
196	aten.atan.complex,	196	aten.atan.complex,
197	aten.atan.default,	197	aten.atan.default,
198	aten.atan.float,	198	aten.atan.float,
199	aten.atan.int,	199	aten.atan.int,
200	aten.atan.out,	200	aten.atan.out,
201	aten.atan2,	201	aten.atan2,
202	aten.atan2.Scalar_Scalar,	202	aten.atan2.Scalar_Scalar,
203	aten.atan2.default,	203	aten.atan2.default,
204	aten.atan2.float,	204	aten.atan2.float,
205	aten.atan2.float_int,	205	aten.atan2.float_int,
206	aten.atan2.int,	206	aten.atan2.int,
207	aten.atan2.int_float,	207	aten.atan2.int_float,
208	aten.atan2.out,	208	aten.atan2.out,
209	aten.atanh,	209	aten.atanh,
210	aten.atanh.Scalar,	210	aten.atanh.Scalar,
211	aten.atanh.complex,	211	aten.atanh.complex,
212	aten.atanh.default,	212	aten.atanh.default,
213	aten.atanh.float,	213	aten.atanh.float,
214	aten.atanh.int,	214	aten.atanh.int,
215	aten.atanh.out,	215	aten.atanh.out,
216	aten.avg_pool2d,	216	aten.avg_pool2d,
217	aten.avg_pool2d.default,	217	aten.avg_pool2d.default,
218	aten.avg_pool2d.out,	218	aten.avg_pool2d.out,
219	aten.avg_pool2d_backward,	219	aten.avg_pool2d_backward,
220	aten.avg_pool2d_backward.default,	220	aten.avg_pool2d_backward.default,
221	aten.avg_pool2d_backward.grad_input,	221	aten.avg_pool2d_backward.grad_input,
222	aten.avg_pool3d,	222	aten.avg_pool3d,
223	aten.avg_pool3d.default,	223	aten.avg_pool3d.default,
224	aten.avg_pool3d.out,	224	aten.avg_pool3d.out,
225	aten.avg_pool3d_backward,	225	aten.avg_pool3d_backward,
226	aten.avg_pool3d_backward.default,	226	aten.avg_pool3d_backward.default,
227	aten.avg_pool3d_backward.grad_input,	227	aten.avg_pool3d_backward.grad_input,
228	aten.baddbmm,	228	aten.baddbmm,
229	aten.baddbmm.default,	229	aten.baddbmm.default,
230	aten.baddbmm.out,	230	aten.baddbmm.out,
231	aten.bernoulli.p,	231	aten.bernoulli.p,
232	aten.bernoulli_,	232	aten.bernoulli_,
233	aten.bernoulli_.Tensor,	233	aten.bernoulli_.Tensor,
234	aten.bernoulli_.float,	234	aten.bernoulli_.float,
235	aten.bitwise_and_,	235	aten.bitwise_and_,
236	aten.bitwise_left_shift,	236	aten.bitwise_left_shift,
237	aten.bitwise_left_shift.Scalar_Tensor,	237	aten.bitwise_left_shift.Scalar_Tensor,
238	aten.bitwise_left_shift.Scalar_Tensor_out,	238	aten.bitwise_left_shift.Scalar_Tensor_out,
239	aten.bitwise_left_shift.Tensor,	239	aten.bitwise_left_shift.Tensor,
240	aten.bitwise_left_shift.Tensor_Scalar,	240	aten.bitwise_left_shift.Tensor_Scalar,
241	aten.bitwise_left_shift.Tensor_Scalar_out,	241	aten.bitwise_left_shift.Tensor_Scalar_out,
242	aten.bitwise_left_shift.Tensor_out,	242	aten.bitwise_left_shift.Tensor_out,
243	aten.bitwise_left_shift_,	243	aten.bitwise_left_shift_,
244	aten.bitwise_left_shift_.Tensor,	244	aten.bitwise_left_shift_.Tensor,
245	aten.bitwise_left_shift_.Tensor_Scalar,	245	aten.bitwise_left_shift_.Tensor_Scalar,
246	aten.bitwise_not,	246	aten.bitwise_not,
247	aten.bitwise_not.default,	247	aten.bitwise_not.default,
248	aten.bitwise_not.out,	248	aten.bitwise_not.out,
249	aten.bitwise_not_,	249	aten.bitwise_not_,
250	aten.bitwise_not_.default,	250	aten.bitwise_not_.default,
251	aten.bitwise_or,	251	aten.bitwise_or,
252	aten.bitwise_or.Scalar,	252	aten.bitwise_or.Scalar,
253	aten.bitwise_or.Tensor,	253	aten.bitwise_or.Tensor,
254	aten.bitwise_or_,	254	aten.bitwise_or_,
255	aten.bitwise_right_shift,	255	aten.bitwise_right_shift,
256	aten.bitwise_right_shift.Scalar_Tensor,	256	aten.bitwise_right_shift.Scalar_Tensor,
257	aten.bitwise_right_shift.Scalar_Tensor_out,	257	aten.bitwise_right_shift.Scalar_Tensor_out,
258	aten.bitwise_right_shift.Tensor,	258	aten.bitwise_right_shift.Tensor,
259	aten.bitwise_right_shift.Tensor_Scalar,	259	aten.bitwise_right_shift.Tensor_Scalar,
260	aten.bitwise_right_shift.Tensor_Scalar_out,	260	aten.bitwise_right_shift.Tensor_Scalar_out,
261	aten.bitwise_right_shift.Tensor_out,	261	aten.bitwise_right_shift.Tensor_out,
262	aten.bitwise_right_shift_,	262	aten.bitwise_right_shift_,
263	aten.bitwise_right_shift_.Tensor,	263	aten.bitwise_right_shift_.Tensor,
264	aten.bitwise_right_shift_.Tensor_Scalar,	264	aten.bitwise_right_shift_.Tensor_Scalar,
265	aten.bitwise_xor_,	265	aten.bitwise_xor_,
266	aten.bitwise_xor_.Scalar,	266	aten.bitwise_xor_.Scalar,
267	aten.bitwise_xor_.Tensor,	267	aten.bitwise_xor_.Tensor,
268	aten.broadcast_tensors,	268	aten.broadcast_tensors,
269	aten.bucketize,	269	aten.bucketize,
270	aten.bucketize.Scalar,	270	aten.bucketize.Scalar,
271	aten.bucketize.Scalar_out,	271	aten.bucketize.Scalar_out,
272	aten.bucketize.Tensor,	272	aten.bucketize.Tensor,
273	aten.bucketize.Tensor_out,	273	aten.bucketize.Tensor_out,
274	aten.cat.names_out,	274	aten.cat.names_out,
275	aten.cat.out,	275	aten.cat.out,
276	aten.constant_pad_nd,	276	aten.constant_pad_nd,
277	aten.constant_pad_nd.default,	277	aten.constant_pad_nd.default,
278	aten.constant_pad_nd.out,	278	aten.constant_pad_nd.out,
279	aten.convolution,	279	aten.convolution,
280	aten.convolution.default,	280	aten.convolution.default,
281	aten.convolution.out,	281	aten.convolution.out,
282	aten.copysign,	282	aten.copysign,
283	aten.copysign.Scalar,	283	aten.copysign.Scalar,
284	aten.copysign.Scalar_out,	284	aten.copysign.Scalar_out,
285	aten.copysign.Tensor,	285	aten.copysign.Tensor,
286	aten.copysign.default,	286	aten.copysign.default,
287	aten.copysign.float,	287	aten.copysign.float,
288	aten.copysign.float_int,	288	aten.copysign.float_int,
289	aten.copysign.int,	289	aten.copysign.int,
290	aten.copysign.int_float,	290	aten.copysign.int_float,
291	aten.copysign.out,	291	aten.copysign.out,
292	aten.cosh,	292	aten.cosh,
293	aten.cosh.Scalar,	293	aten.cosh.Scalar,
294	aten.cosh.complex,	294	aten.cosh.complex,
295	aten.cosh.default,	295	aten.cosh.default,
296	aten.cosh.float,	296	aten.cosh.float,
297	aten.cosh.int,	297	aten.cosh.int,
298	aten.cosh.out,	298	aten.cosh.out,
299	aten.cummax,	299	aten.cummax,
300	aten.cummax.dimname_out,	300	aten.cummax.dimname_out,
301	aten.cummax.out,	301	aten.cummax.out,
302	aten.cummin,	302	aten.cummin,
303	aten.cummin.dimname_out,	303	aten.cummin.dimname_out,
304	aten.cummin.out,	304	aten.cummin.out,
305	aten.cumprod,	305	aten.cumprod,
306	aten.cumprod.dimname_out,	306	aten.cumprod.dimname_out,
307	aten.cumprod.out,	307	aten.cumprod.out,
308	aten.cumsum.dimname_out,	308	aten.cumsum.dimname_out,
309	aten.cumsum.out,	309	aten.cumsum.out,
310	aten.detach,	310	aten.detach,
311	aten.detach_,	311	aten.detach_,
312	aten.diagonal,	312	aten.diagonal,
313	aten.diagonal.Dimname,	313	aten.diagonal.Dimname,
314	aten.diagonal.default,	314	aten.diagonal.default,
315	aten.diagonal_copy,	315	aten.diagonal_copy,
316	aten.diagonal_scatter,	316	aten.diagonal_scatter,
317	aten.diagonal_scatter.default,	317	aten.diagonal_scatter.default,
318	aten.diagonal_scatter.out,	318	aten.diagonal_scatter.out,
319	aten.div_.Tensor,	319	aten.div_.Tensor,
320	aten.div_.Tensor_mode,	320	aten.div_.Tensor_mode,
321	aten.empty,	321	aten.empty,
322	aten.empty.memory_format,	322	aten.empty.memory_format,
323	aten.empty.names,	323	aten.empty.names,
324	aten.empty.names_out,	324	aten.empty.names_out,
325	aten.empty.out,	325	aten.empty.out,
326	aten.empty_like,	326	aten.empty_like,
327	aten.empty_strided,	327	aten.empty_strided,
328	aten.empty_strided.default,	328	aten.empty_strided.default,
329	aten.empty_strided.out,	329	aten.empty_strided.out,
330	aten.erfc,	330	aten.erfc,
331	aten.erfc.Scalar,	331	aten.erfc.Scalar,
332	aten.erfc.default,	332	aten.erfc.default,
333	aten.erfc.float,	333	aten.erfc.float,
334	aten.erfc.int,	334	aten.erfc.int,
335	aten.erfc.out,	335	aten.erfc.out,
336	aten.erfinv,	336	aten.erfinv,
337	aten.erfinv.default,	337	aten.erfinv.default,
338	aten.erfinv.out,	338	aten.erfinv.out,
339	aten.exp2.out,	339	aten.exp2.out,
340	aten.expand_as,	340	aten.expand_as,
341	aten.expand_as.default,	341	aten.expand_as.default,
342	aten.expm1,	342	aten.expm1,
343	aten.expm1.Scalar,	343	aten.expm1.Scalar,
344	aten.expm1.default,	344	aten.expm1.default,
345	aten.expm1.float,	345	aten.expm1.float,
346	aten.expm1.int,	346	aten.expm1.int,
347	aten.expm1.out,	347	aten.expm1.out,
348	aten.fill_,	348	aten.fill_,
349	aten.fmod,	349	aten.fmod,
350	aten.fmod.Scalar,	350	aten.fmod.Scalar,
351	aten.fmod.Scalar_out,	351	aten.fmod.Scalar_out,
352	aten.fmod.Tensor,	352	aten.fmod.Tensor,
353	aten.fmod.Tensor_out,	353	aten.fmod.Tensor_out,
354	aten.fmod.default,	354	aten.fmod.default,
355	aten.fmod.float,	355	aten.fmod.float,
356	aten.fmod.float_int,	356	aten.fmod.float_int,
357	aten.fmod.int,	357	aten.fmod.int,
358	aten.fmod.int_float,	358	aten.fmod.int_float,
359	aten.fractional_max_pool2d,	359	aten.fractional_max_pool2d,
360	aten.fractional_max_pool2d.default,	360	aten.fractional_max_pool2d.default,
361	aten.fractional_max_pool2d.output,	361	aten.fractional_max_pool2d.output,
362	aten.frexp,	362	aten.frexp,
363	aten.frexp.Tensor,	363	aten.frexp.Tensor,
364	aten.frexp.Tensor_out,	364	aten.frexp.Tensor_out,
365	aten.frexp.default,	365	aten.frexp.default,
366	aten.full_like,	366	aten.full_like,
367	aten.glu,	367	aten.glu,
368	aten.glu.default,	368	aten.glu.default,
369	aten.hypot,	369	aten.hypot,
370	aten.hypot.default,	370	aten.hypot.default,
371	aten.hypot.out,	371	aten.hypot.out,
372	aten.i0.default,	372	aten.i0.default,
373	aten.i0.out,	373	aten.i0.out,
374	aten.index.Tensor_out,	374	aten.index.Tensor_out,
375	aten.index.list_Tensor,	375	aten.index.list_Tensor,
376	aten.index.list_bool,	376	aten.index.list_bool,
377	aten.index.list_float,	377	aten.index.list_float,
378	aten.index.list_int,	378	aten.index.list_int,
379	aten.index.list_str,	379	aten.index.list_str,
380	aten.index.str,	380	aten.index.str,
381	aten.index_put.hacked_twin,	381	aten.index_put.hacked_twin,
382	aten.index_put_.hacked_twin,	382	aten.index_put_.hacked_twin,
383	aten.isinf,	383	aten.isinf,
384	aten.isinf.complex,	384	aten.isinf.complex,
385	aten.isinf.default,	385	aten.isinf.default,
386	aten.isinf.float,	386	aten.isinf.float,
387	aten.isinf.out,	387	aten.isinf.out,
388	aten.lgamma,	388	aten.lgamma,
389	aten.lgamma.Scalar,	389	aten.lgamma.Scalar,
390	aten.lgamma.default,	390	aten.lgamma.default,
391	aten.lgamma.float,	391	aten.lgamma.float,
392	aten.lgamma.int,	392	aten.lgamma.int,
393	aten.lgamma.out,	393	aten.lgamma.out,
394	aten.lift,	394	aten.lift,
395	aten.lift_fresh,	395	aten.lift_fresh,
396	aten.lift_fresh.default,	396	aten.lift_fresh.default,
397	aten.lift_fresh_copy,	397	aten.lift_fresh_copy,
398	aten.lift_fresh_copy.default,	398	aten.lift_fresh_copy.default,
399	aten.lift_fresh_copy.out,	399	aten.lift_fresh_copy.out,
400	aten.log10,	400	aten.log10,
401	aten.log10.Scalar,	401	aten.log10.Scalar,
402	aten.log10.complex,	402	aten.log10.complex,
403	aten.log10.default,	403	aten.log10.default,
404	aten.log10.float,	404	aten.log10.float,
405	aten.log10.int,	405	aten.log10.int,
406	aten.log10.out,	406	aten.log10.out,
407	aten.logcumsumexp,	407	aten.logcumsumexp,
408	aten.logcumsumexp.dimname_out,	408	aten.logcumsumexp.dimname_out,
409	aten.logcumsumexp.out,	409	aten.logcumsumexp.out,
410	aten.logical_and_,	410	aten.logical_and_,
411	aten.logical_and_.default,	411	aten.logical_and_.default,
412	aten.logical_not_,	412	aten.logical_not_,
413	aten.logical_not_.default,	413	aten.logical_not_.default,
414	aten.logical_or_,	414	aten.logical_or_,
415	aten.logical_or_.default,	415	aten.logical_or_.default,
416	aten.logical_xor,	416	aten.logical_xor,
417	aten.logical_xor.default,	417	aten.logical_xor.default,
418	aten.logical_xor.out,	418	aten.logical_xor.out,
419	aten.logical_xor_,	419	aten.logical_xor_,
420	aten.logical_xor_.default,	420	aten.logical_xor_.default,
421	aten.max_pool2d_with_indices,	421	aten.max_pool2d_with_indices,
422	aten.max_pool2d_with_indices_backward,	422	aten.max_pool2d_with_indices_backward,
423	aten.max_pool2d_with_indices_backward.default,	423	aten.max_pool2d_with_indices_backward.default,
424	aten.max_pool2d_with_indices_backward.grad_input,	424	aten.max_pool2d_with_indices_backward.grad_input,
425	aten.mul_,	425	aten.mul_,
426	aten.mul_.Scalar,	426	aten.mul_.Scalar,
427	aten.mul_.Tensor,	427	aten.mul_.Tensor,
428	aten.mul_.t,	428	aten.mul_.t,
429	aten.native_dropout,	429	aten.native_dropout,
430	aten.native_dropout.default,	430	aten.native_dropout.default,
431	aten.native_dropout.out,	431	aten.native_dropout.out,
432	aten.new_empty,	432	aten.new_empty,
433	aten.new_empty_strided,	433	aten.new_empty_strided,
434	aten.new_empty_strided.default,	434	aten.new_empty_strided.default,
435	aten.new_empty_strided.out,	435	aten.new_empty_strided.out,
436	aten.nextafter,	436	aten.nextafter,
437	aten.nextafter.default,	437	aten.nextafter.default,
438	aten.nextafter.out,	438	aten.nextafter.out,
439	aten.prod,	439	aten.prod,
440	aten.prod.Dimname_out,	440	aten.prod.Dimname_out,
441	aten.prod.default,	441	aten.prod.default,
442	aten.prod.dim_int,	442	aten.prod.dim_int,
443	aten.prod.int_out,	443	aten.prod.int_out,
444	aten.prod.out,	444	aten.prod.out,
445	aten.rand,	445	aten.rand,
446	aten.rand.generator_with_names,	446	aten.rand.generator_with_names,
447	aten.rand.generator_with_names_out,	447	aten.rand.generator_with_names_out,
448	aten.rand.names,	448	aten.rand.names,
449	aten.rand.names_out,	449	aten.rand.names_out,
450	aten.rand.out,	450	aten.rand.out,
451	aten.randn,	451	aten.randn,
452	aten.randn.generator_with_names,	452	aten.randn.generator_with_names,
453	aten.randn.generator_with_names_out,	453	aten.randn.generator_with_names_out,
454	aten.randn.names,	454	aten.randn.names,
455	aten.randn.names_out,	455	aten.randn.names_out,
456	aten.randn.out,	456	aten.randn.out,
457	aten.relu_,	457	aten.relu_,
458	aten.relu_.default,	458	aten.relu_.default,
459	aten.remainder,	459	aten.remainder,
460	aten.remainder.Scalar,	460	aten.remainder.Scalar,
461	aten.remainder.Scalar_Tensor_out,	461	aten.remainder.Scalar_Tensor_out,
462	aten.remainder.Scalar_out,	462	aten.remainder.Scalar_out,
463	aten.remainder.Tensor,	463	aten.remainder.Tensor,
464	aten.remainder.Tensor_out,	464	aten.remainder.Tensor_out,
465	aten.remainder.default,	465	aten.remainder.default,
466	aten.remainder.float,	466	aten.remainder.float,
467	aten.remainder.float_int,	467	aten.remainder.float_int,
468	aten.remainder.int,	468	aten.remainder.int,
469	aten.remainder.int_float,	469	aten.remainder.int_float,
470	aten.resize,	470	aten.resize,
471	aten.resize.default,	471	aten.resize.default,
472	aten.resize.out,	472	aten.resize.out,
473	aten.round.default,	473	aten.round.default,
474	aten.scatter.reduce,	474	aten.scatter.reduce,
475	aten.scatter.reduce_out,	475	aten.scatter.reduce_out,
476	aten.scatter.src_out,	476	aten.scatter.src_out,
477	aten.scatter.value,	477	aten.scatter.value,
478	aten.scatter.value_out,	478	aten.scatter.value_out,
479	aten.scatter.value_reduce,	479	aten.scatter.value_reduce,
480	aten.scatter.value_reduce_out,	480	aten.scatter.value_reduce_out,
481	aten.scatter_.reduce,	481	aten.scatter_.reduce,
482	aten.scatter_.value,	482	aten.scatter_.value,
483	aten.scatter_.value_reduce,	483	aten.scatter_.value_reduce,
484	aten.scatter_add,	484	aten.scatter_add,
485	aten.scatter_add.default,	485	aten.scatter_add.default,
486	aten.scatter_add.out,	486	aten.scatter_add.out,
487	aten.scatter_add_,	487	aten.scatter_add_,
488	aten.scatter_add_.default,	488	aten.scatter_add_.default,
489	aten.scatter_reduce.two,	489	aten.scatter_reduce.two,
490	aten.scatter_reduce.two_out,	490	aten.scatter_reduce.two_out,
491	aten.scatter_reduce_.two,	491	aten.scatter_reduce_.two,
492	aten.searchsorted.Tensor,	492	aten.searchsorted.Tensor,
493	aten.set_.source_Tensor,	493	aten.set_.source_Tensor,
494	aten.sigmoid_,	494	aten.sigmoid_,
495	aten.sigmoid_.default,	495	aten.sigmoid_.default,
496	aten.signbit,	496	aten.signbit,
497	aten.signbit.default,	497	aten.signbit.default,
498	aten.signbit.out,	498	aten.signbit.out,
499	aten.sinh,	499	aten.sinh,
500	aten.sinh.Scalar,	500	aten.sinh.Scalar,
501	aten.sinh.complex,	501	aten.sinh.complex,
502	aten.sinh.default,	502	aten.sinh.default,
503	aten.sinh.float,	503	aten.sinh.float,
504	aten.sinh.int,	504	aten.sinh.int,
505	aten.sinh.out,	505	aten.sinh.out,
506	aten.special_bessel_j0.default,	506	aten.special_bessel_j0.default,
507	aten.special_bessel_j0.out,	507	aten.special_bessel_j0.out,
508	aten.special_bessel_j1.default,	508	aten.special_bessel_j1.default,
509	aten.special_bessel_j1.out,	509	aten.special_bessel_j1.out,
510	aten.special_bessel_y0.default,	510	aten.special_bessel_y0.default,
511	aten.special_bessel_y0.out,	511	aten.special_bessel_y0.out,
512	aten.special_bessel_y1.default,	512	aten.special_bessel_y1.default,
513	aten.special_bessel_y1.out,	513	aten.special_bessel_y1.out,
514	aten.special_erf,	514	aten.special_erf,
515	aten.special_erf.out,	515	aten.special_erf.out,
516	aten.special_erfcx.default,	516	aten.special_erfcx.default,
517	aten.special_erfcx.out,	517	aten.special_erfcx.out,
518	aten.special_i1.default,	518	aten.special_i1.default,
519	aten.special_i1.out,	519	aten.special_i1.out,
520	aten.special_modified_bessel_i0.default,	520	aten.special_modified_bessel_i0.default,
521	aten.special_modified_bessel_i0.out,	521	aten.special_modified_bessel_i0.out,
522	aten.special_modified_bessel_i1.default,	522	aten.special_modified_bessel_i1.default,
523	aten.special_modified_bessel_i1.out,	523	aten.special_modified_bessel_i1.out,
524	aten.split.str,	524	aten.split.str,
525	aten.square,	525	aten.square,
526	aten.square.out,	526	aten.square.out,
527	aten.squeeze_,	527	aten.squeeze_,
528	aten.squeeze_.default,	528	aten.squeeze_.default,
529	aten.squeeze_.dim,	529	aten.squeeze_.dim,
530	aten.squeeze_.dimname,	530	aten.squeeze_.dimname,
531	aten.squeeze_.dims,	531	aten.squeeze_.dims,
532	aten.squeeze_copy,	532	aten.squeeze_copy,
533	aten.sub_,	533	aten.sub_,
534	aten.sub_.Scalar,	534	aten.sub_.Scalar,
535	aten.sub_.Tensor,	535	aten.sub_.Tensor,
536	aten.sym_constrain_range,	536	aten.sym_constrain_range,
537	aten.sym_constrain_range.default,	537	aten.sym_constrain_range.default,
538	aten.sym_numel,	538	aten.sym_numel,
539	aten.sym_numel.default,	539	aten.sym_numel.default,
540	aten.tan,	540	aten.tan,
541	aten.tan.Scalar,	541	aten.tan.Scalar,
542	aten.tan.complex,	542	aten.tan.complex,
543	aten.tan.default,	543	aten.tan.default,
544	aten.tan.float,	544	aten.tan.float,
545	aten.tan.int,	545	aten.tan.int,
546	aten.tan.out,	546	aten.tan.out,
547	aten.true_divide,	547	aten.true_divide,
548	aten.true_divide.out,	548	aten.true_divide.out,
549	aten.trunc,	549	aten.trunc,
550	aten.trunc.default,	550	aten.trunc.default,
551	aten.trunc.out,	551	aten.trunc.out,
552	aten.unbind,	552	aten.unbind,
553	aten.unbind.Dimname,	553	aten.unbind.Dimname,
554	aten.unbind.int,	554	aten.unbind.int,
555	aten.unfold,	555	aten.unfold,
556	aten.unfold.default,	556	aten.unfold.default,
557	aten.unsqueeze_,	557	aten.unsqueeze_,
558	aten.unsqueeze_.default,	558	aten.unsqueeze_.default,
559	aten.upsample_nearest2d_backward.default,	559	aten.upsample_nearest2d_backward.default,
560	aten.var.correction,	560	aten.var.correction,
561	aten.var.correction_names_out,	561	aten.var.correction_names_out,
562	aten.var.correction_out,	562	aten.var.correction_out,
563	aten.var.names_out,	563	aten.var.names_out,
564	aten.var.out,	564	aten.var.out,
565	aten.var_mean.correction,	565	aten.var_mean.correction,
566	aten.var_mean.correction_out,	566	aten.var_mean.correction_out,
567	aten.view,	567	aten.view,
568	aten.view.default,	568	aten.view.default,
569	aten.view.dtype,	569	aten.view.dtype,
570	cond,	570	cond,
571	fsdp.copy_.default,	571	fsdp.copy_.default,
572	inductor.resize_storage_bytes_,	572	inductor.resize_storage_bytes_,
573	inductor.resize_storage_bytes_.default,	573	inductor.resize_storage_bytes_.default,
574	invoke_quant,	574	invoke_quant,
575	prims._low_memory_max_pool2d_offsets_to_indices,	575	prims._low_memory_max_pool2d_offsets_to_indices,
576	prims._low_memory_max_pool2d_offsets_to_indices.default,	576	prims._low_memory_max_pool2d_offsets_to_indices.default,
577	prims._low_memory_max_pool2d_with_offsets,	577	prims._low_memory_max_pool2d_with_offsets,
578	prims._low_memory_max_pool2d_with_offsets.default,	578	prims._low_memory_max_pool2d_with_offsets.default,
579	prims._sink_tokens.default,	579	prims._sink_tokens.default,
580	prims._unsafe_index_put_.default,	580	prims._unsafe_index_put_.default,
581	prims.abs,	581	prims.abs,
582	prims.abs.default,	582	prims.abs.default,
583	prims.acos,	583	prims.acos,
584	prims.acos.default,	584	prims.acos.default,
585	prims.acosh,	585	prims.acosh,
586	prims.acosh.default,	586	prims.acosh.default,
587	prims.add,	587	prims.add,
588	prims.add.default,	588	prims.add.default,
589	prims.asin,	589	prims.asin,
590	prims.asin.default,	590	prims.asin.default,
591	prims.asinh,	591	prims.asinh,
592	prims.asinh.default,	592	prims.asinh.default,
593	prims.atan,	593	prims.atan,
594	prims.atan.default,	594	prims.atan.default,
595	prims.atan2,	595	prims.atan2,
596	prims.atan2.default,	596	prims.atan2.default,
597	prims.atanh,	597	prims.atanh,
598	prims.atanh.default,	598	prims.atanh.default,
599	prims.bessel_j0,	599	prims.bessel_j0,
600	prims.bessel_j0.default,	600	prims.bessel_j0.default,
601	prims.bessel_j1,	601	prims.bessel_j1,
602	prims.bessel_j1.default,	602	prims.bessel_j1.default,
603	prims.bitwise_and,	603	prims.bitwise_and,
604	prims.bitwise_and.default,	604	prims.bitwise_and.default,
605	prims.bitwise_not,	605	prims.bitwise_not,
606	prims.bitwise_not.default,	606	prims.bitwise_not.default,
607	prims.bitwise_or,	607	prims.bitwise_or,
608	prims.bitwise_or.default,	608	prims.bitwise_or.default,
609	prims.bitwise_xor,	609	prims.bitwise_xor,
610	prims.bitwise_xor.default,	610	prims.bitwise_xor.default,
611	prims.ceil,	611	prims.ceil,
612	prims.ceil.default,	612	prims.ceil.default,
613	prims.cos,	613	prims.cos,
614	prims.cos.default,	614	prims.cos.default,
615	prims.cosh,	615	prims.cosh,
616	prims.cosh.default,	616	prims.cosh.default,
617	prims.digamma,	617	prims.digamma,
618	prims.div,	618	prims.div,
619	prims.div.default,	619	prims.div.default,
620	prims.eq,	620	prims.eq,
621	prims.eq.default,	621	prims.eq.default,
622	prims.erf,	622	prims.erf,
623	prims.erf.default,	623	prims.erf.default,
624	prims.erfc,	624	prims.erfc,
625	prims.erfc.default,	625	prims.erfc.default,
626	prims.erfcx,	626	prims.erfcx,
627	prims.erfcx.default,	627	prims.erfcx.default,
628	prims.exp,	628	prims.exp,
629	prims.exp.default,	629	prims.exp.default,
630	prims.expm1,	630	prims.expm1,
631	prims.expm1.default,	631	prims.expm1.default,
632	prims.fma,	632	prims.fma,
633	prims.fma.default,	633	prims.fma.default,
634	prims.fmod,	634	prims.fmod,
635	prims.fmod.default,	635	prims.fmod.default,
636	prims.frexp,	636	prims.frexp,
637	prims.frexp.default,	637	prims.frexp.default,
638	prims.ge,	638	prims.ge,
639	prims.ge.default,	639	prims.ge.default,
640	prims.gt,	640	prims.gt,
641	prims.gt.default,	641	prims.gt.default,
642	prims.hypot,	642	prims.hypot,
643	prims.hypot.default,	643	prims.hypot.default,
644	prims.igamma,	644	prims.igamma,
645	prims.igammac,	645	prims.igammac,
646	prims.inductor_force_stride_order.default,	646	prims.inductor_force_stride_order.default,
647	prims.inductor_lookup_seed.default,	647	prims.inductor_lookup_seed.default,
648	prims.inductor_randint.default,	648	prims.inductor_randint.default,
649	prims.inductor_random.default,	649	prims.inductor_random.default,
650	prims.inductor_seed.default,	650	prims.inductor_seed.default,
651	prims.inductor_seeds.default,	651	prims.inductor_seeds.default,
652	prims.le,	652	prims.le,
653	prims.le.default,	653	prims.le.default,
654	prims.lgamma,	654	prims.lgamma,
655	prims.lgamma.default,	655	prims.lgamma.default,
656	prims.log,	656	prims.log,
657	prims.log.default,	657	prims.log.default,
658	prims.log10,	658	prims.log10,
659	prims.log10.default,	659	prims.log10.default,
660	prims.log1p,	660	prims.log1p,
661	prims.log1p.default,	661	prims.log1p.default,
662	prims.lt,	662	prims.lt,
663	prims.lt.default,	663	prims.lt.default,
664	prims.minimum,	664	prims.minimum,
665	prims.minimum.default,	665	prims.minimum.default,
666	prims.ndtri,	666	prims.ndtri,
667	prims.ndtri.default,	667	prims.ndtri.default,
668	prims.ne,	668	prims.ne,
669	prims.ne.default,	669	prims.ne.default,
670	prims.neg,	670	prims.neg,
671	prims.neg.default,	671	prims.neg.default,
672	prims.nextafter,	672	prims.nextafter,
673	prims.nextafter.default,	673	prims.nextafter.default,
674	prims.prepare_softmax_online.default,	674	prims.prepare_softmax_online.default,
675	prims.reciprocal,	675	prims.reciprocal,
676	prims.reciprocal.default,	676	prims.reciprocal.default,
677	prims.remainder,	677	prims.remainder,
678	prims.remainder.default,	678	prims.remainder.default,
679	prims.rev.default,	679	prims.rev.default,
680	prims.rsqrt.default,	680	prims.rsqrt.default,
681	prims.sign,	681	prims.sign,
682	prims.sign.default,	682	prims.sign.default,
683	prims.signbit,	683	prims.signbit,
684	prims.signbit.default,	684	prims.signbit.default,
685	prims.sin,	685	prims.sin,
686	prims.sin.default,	686	prims.sin.default,
687	prims.sinh,	687	prims.sinh,
688	prims.sinh.default,	688	prims.sinh.default,
689	prims.spherical_bessel_j0,	689	prims.spherical_bessel_j0,
690	prims.spherical_bessel_j0.default,	690	prims.spherical_bessel_j0.default,
691	prims.sqrt,	691	prims.sqrt,
692	prims.sqrt.default,	692	prims.sqrt.default,
693	prims.sub,	693	prims.sub,
694	prims.sub.default,	694	prims.sub.default,
695	prims.sum,	695	prims.sum,
696	prims.sum.default,	696	prims.sum.default,
697	prims.tan,	697	prims.tan,
698	prims.tan.default,	698	prims.tan.default,
699	prims.var,	699	prims.var,
700	prims.var.default,	700	prims.var.default,
701	prims.view_of,	701	prims.view_of,
702	prims.view_of.default,	702	prims.view_of.default,
703	prims.xor_sum,	703	prims.xor_sum,
704	prims.xor_sum.default,	704	prims.xor_sum.default,
705	prims.zeta,	705	prims.zeta,
706	prims.zeta.default,	706	prims.zeta.default,
707	quantized_decomposed.dequantize_per_channel,	707	quantized_decomposed.dequantize_per_channel,
708	quantized_decomposed.dequantize_per_channel.default,	708	quantized_decomposed.dequantize_per_channel.default,
709	quantized_decomposed.dequantize_per_tensor.default,	709	quantized_decomposed.dequantize_per_tensor.default,
710	quantized_decomposed.dequantize_per_tensor.tensor,	710	quantized_decomposed.dequantize_per_tensor.tensor,
711	quantized_decomposed.quantize_per_channel,	711	quantized_decomposed.quantize_per_channel,
712	quantized_decomposed.quantize_per_channel.default,	712	quantized_decomposed.quantize_per_channel.default,
713	quantized_decomposed.quantize_per_tensor.default,	713	quantized_decomposed.quantize_per_tensor.default,
714	quantized_decomposed.quantize_per_tensor.tensor,	714	quantized_decomposed.quantize_per_tensor.tensor,
715	rngprims.philox_rand,	715	rngprims.philox_rand,
716	rngprims.philox_rand.default,	716	rngprims.philox_rand.default,
717	while_loop,	717	while_loop,
718	with_effects,	718	with_effects,
719	]	719	]
720	TORCH_NATIVE_FALLBACK_LIST = [	720	TORCH_NATIVE_FALLBACK_LIST = [
721	_quantized.wrapped_fbgemm_linear_fp16_weight.default,	721	_quantized.wrapped_fbgemm_linear_fp16_weight.default,
722	_quantized.wrapped_fbgemm_pack_gemm_matrix_fp16.default,	722	_quantized.wrapped_fbgemm_pack_gemm_matrix_fp16.default,
723	aten._adaptive_avg_pool2d_backward.default,	723	aten._adaptive_avg_pool2d_backward.default,
724	aten._adaptive_avg_pool2d_backward.out,	724	aten._adaptive_avg_pool2d_backward.out,
725	aten._adaptive_avg_pool3d.default,	725	aten._adaptive_avg_pool3d.default,
726	aten._adaptive_avg_pool3d.out,	726	aten._adaptive_avg_pool3d.out,
727	aten._adaptive_avg_pool3d_backward.default,	727	aten._adaptive_avg_pool3d_backward.default,
728	aten._adaptive_avg_pool3d_backward.out,	728	aten._adaptive_avg_pool3d_backward.out,
729	aten._addmm_activation.default,	729	aten._addmm_activation.default,
730	aten._addmm_activation.out,	730	aten._addmm_activation.out,
731	aten._cdist_backward.default,	731	aten._cdist_backward.default,
732	aten._cdist_backward.out,	732	aten._cdist_backward.out,
733	aten._cdist_forward.default,	733	aten._cdist_forward.default,
734	aten._cdist_forward.out,	734	aten._cdist_forward.out,
735	aten._cudnn_rnn.default,	735	aten._cudnn_rnn.default,
736	aten._cudnn_rnn.out,	736	aten._cudnn_rnn.out,
737	aten._cudnn_rnn_backward.default,	737	aten._cudnn_rnn_backward.default,
738	aten._cudnn_rnn_backward.out,	738	aten._cudnn_rnn_backward.out,
739	aten._dyn_quant_matmul_4bit.default,	739	aten._dyn_quant_matmul_4bit.default,
740	aten._dyn_quant_pack_4bit_weight.default,	740	aten._dyn_quant_pack_4bit_weight.default,
741	aten._efficient_attention_backward.default,	741	aten._efficient_attention_backward.default,
742	aten._efficient_attention_forward.default,	742	aten._efficient_attention_forward.default,
743	aten._efficientzerotensor.default,	743	aten._efficientzerotensor.default,
744	aten._efficientzerotensor.out,	744	aten._efficientzerotensor.out,
745	aten._embedding_bag.default,	745	aten._embedding_bag.default,
746	aten._embedding_bag.out,	746	aten._embedding_bag.out,
747	aten._embedding_bag_backward.default,	747	aten._embedding_bag_backward.default,
748	aten._embedding_bag_forward_only.default,	748	aten._embedding_bag_forward_only.default,
749	aten._embedding_bag_forward_only.out,	749	aten._embedding_bag_forward_only.out,
750	aten._embedding_bag_per_sample_weights_backward.default,	750	aten._embedding_bag_per_sample_weights_backward.default,
751	aten._embedding_bag_per_sample_weights_backward.out,	751	aten._embedding_bag_per_sample_weights_backward.out,
752	aten._fft_r2c.default,	752	aten._fft_r2c.default,
753	aten._fft_r2c.out,	753	aten._fft_r2c.out,
754	aten._flash_attention_backward.default,	754	aten._flash_attention_backward.default,
755	aten._flash_attention_forward.default,	755	aten._flash_attention_forward.default,
756	aten._fused_moving_avg_obs_fq_helper.default,	756	aten._fused_moving_avg_obs_fq_helper.default,
757	aten._fused_moving_avg_obs_fq_helper.out,	757	aten._fused_moving_avg_obs_fq_helper.out,
758	aten._fused_moving_avg_obs_fq_helper_functional.default,	758	aten._fused_moving_avg_obs_fq_helper_functional.default,
759	aten._histogramdd_bin_edges.default,	759	aten._histogramdd_bin_edges.default,
760	aten._histogramdd_from_bin_cts.default,	760	aten._histogramdd_from_bin_cts.default,
761	aten._linalg_check_errors.default,	761	aten._linalg_check_errors.default,
762	aten._linalg_det.default,	762	aten._linalg_det.default,
763	aten._linalg_det.result,	763	aten._linalg_det.result,
764	aten._linalg_eigh.default,	764	aten._linalg_eigh.default,
765	aten._linalg_eigh.eigenvalues,	765	aten._linalg_eigh.eigenvalues,
766	aten._linalg_slogdet.default,	766	aten._linalg_slogdet.default,
767	aten._linalg_slogdet.sign,	767	aten._linalg_slogdet.sign,
768	aten._linalg_solve_ex.default,	768	aten._linalg_solve_ex.default,
769	aten._linalg_solve_ex.result,	769	aten._linalg_solve_ex.result,
770	aten._linalg_svd.U,	770	aten._linalg_svd.U,
771	aten._linalg_svd.default,	771	aten._linalg_svd.default,
772	aten._pdist_backward.default,	772	aten._pdist_backward.default,
773	aten._pdist_backward.out,	773	aten._pdist_backward.out,
774	aten._pdist_forward.default,	774	aten._pdist_forward.default,
775	aten._pdist_forward.out,	775	aten._pdist_forward.out,
776	aten._scaled_dot_product_cudnn_attention.default,	776	aten._scaled_dot_product_cudnn_attention.default,
777	aten._scaled_dot_product_cudnn_attention_backward.default,	777	aten._scaled_dot_product_cudnn_attention_backward.default,
778	aten._scaled_dot_product_efficient_attention.default,	778	aten._scaled_dot_product_efficient_attention.default,
779	aten._scaled_dot_product_efficient_attention_backward.default,	779	aten._scaled_dot_product_efficient_attention_backward.default,
780	aten._scaled_dot_product_flash_attention.default,	780	aten._scaled_dot_product_flash_attention.default,
781	aten._scaled_dot_product_flash_attention_backward.default,	781	aten._scaled_dot_product_flash_attention_backward.default,
782	aten._scaled_dot_product_flash_attention_for_cpu.default,	782	aten._scaled_dot_product_flash_attention_for_cpu.default,
783	aten._scaled_dot_product_flash_attention_for_cpu_backward.default,	783	aten._scaled_dot_product_flash_attention_for_cpu_backward.default,
784	aten._scaled_dot_product_fused_attention_overrideable.default,	784	aten._scaled_dot_product_fused_attention_overrideable.default,
785	aten._scaled_dot_product_fused_attention_overrideable_backward.default,	785	aten._scaled_dot_product_fused_attention_overrideable_backward.default,
786	aten._segment_reduce_backward.default,	786	aten._segment_reduce_backward.default,
787	aten._sparse_coo_tensor_with_dims_and_tensors.default,	787	aten._sparse_coo_tensor_with_dims_and_tensors.default,
788	aten._sparse_coo_tensor_with_dims_and_tensors.out,	788	aten._sparse_coo_tensor_with_dims_and_tensors.out,
789	aten._thnn_fused_lstm_cell.default,	789	aten._thnn_fused_lstm_cell.default,
790	aten._thnn_fused_lstm_cell.out,	790	aten._thnn_fused_lstm_cell.out,
791	aten._to_sparse.default,	791	aten._to_sparse.default,
792	aten._to_sparse.out,	792	aten._to_sparse.out,
793	aten._to_sparse.sparse_dim,	793	aten._to_sparse.sparse_dim,
794	aten._to_sparse.sparse_dim_out,	794	aten._to_sparse.sparse_dim_out,
795	aten._trilinear.default,	795	aten._trilinear.default,
796	aten._trilinear.out,	796	aten._trilinear.out,
797	aten.adaptive_max_pool2d_backward.default,	797	aten.adaptive_max_pool2d_backward.default,
798	aten.adaptive_max_pool2d_backward.grad_input,	798	aten.adaptive_max_pool2d_backward.grad_input,
799	aten.adaptive_max_pool3d.default,	799	aten.adaptive_max_pool3d.default,
800	aten.adaptive_max_pool3d.out,	800	aten.adaptive_max_pool3d.out,
801	aten.adaptive_max_pool3d_backward.default,	801	aten.adaptive_max_pool3d_backward.default,
802	aten.adaptive_max_pool3d_backward.grad_input,	802	aten.adaptive_max_pool3d_backward.grad_input,
803	aten.addbmm.default,	803	aten.addbmm.default,
804	aten.addbmm.out,	804	aten.addbmm.out,
805	aten.angle.Scalar,	805	aten.angle.Scalar,
806	aten.angle.complex,	806	aten.angle.complex,
807	aten.angle.default,	807	aten.angle.default,
808	aten.angle.float,	808	aten.angle.float,
809	aten.angle.int,	809	aten.angle.int,
810	aten.angle.out,	810	aten.angle.out,
811	aten.cholesky_inverse.default,	811	aten.cholesky_inverse.default,
812	aten.cholesky_inverse.out,	812	aten.cholesky_inverse.out,
813	aten.cholesky_solve.default,	813	aten.cholesky_solve.default,
814	aten.cholesky_solve.out,	814	aten.cholesky_solve.out,
815	aten.convolution_backward.default,	815	aten.convolution_backward.default,
816	aten.convolution_backward.out,	816	aten.convolution_backward.out,
817	aten.cummax.default,	817	aten.cummax.default,
818	aten.cummin.default,	818	aten.cummin.default,
819	aten.cumprod.default,	819	aten.cumprod.default,
820	aten.cumsum.default,	820	aten.cumsum.default,
821	aten.digamma.default,	821	aten.digamma.default,
822	aten.digamma.out,	822	aten.digamma.out,
823	aten.exponential.default,	823	aten.exponential.default,
824	aten.fractional_max_pool2d_backward.default,	824	aten.fractional_max_pool2d_backward.default,
825	aten.fractional_max_pool2d_backward.grad_input,	825	aten.fractional_max_pool2d_backward.grad_input,
826	aten.fractional_max_pool3d.default,	826	aten.fractional_max_pool3d.default,
827	aten.fractional_max_pool3d.output,	827	aten.fractional_max_pool3d.output,
828	aten.fractional_max_pool3d_backward.default,	828	aten.fractional_max_pool3d_backward.default,
829	aten.fractional_max_pool3d_backward.grad_input,	829	aten.fractional_max_pool3d_backward.grad_input,
830	aten.gcd.default,	830	aten.gcd.default,
831	aten.geqrf.a,	831	aten.geqrf.a,
832	aten.geqrf.default,	832	aten.geqrf.default,
833	aten.grid_sampler_2d_backward.default,	833	aten.grid_sampler_2d_backward.default,
834	aten.grid_sampler_2d_backward.out,	834	aten.grid_sampler_2d_backward.out,
835	aten.histc.default,	835	aten.histc.default,
836	aten.histc.out,	836	aten.histc.out,
837	aten.histogram.bin_ct,	837	aten.histogram.bin_ct,
838	aten.igamma.default,	838	aten.igamma.default,
839	aten.igamma.out,	839	aten.igamma.out,
840	aten.igammac.default,	840	aten.igammac.default,
841	aten.igammac.out,	841	aten.igammac.out,
842	aten.index_reduce.default,	842	aten.index_reduce.default,
843	aten.index_reduce.out,	843	aten.index_reduce.out,
844	aten.kthvalue.default,	844	aten.kthvalue.default,
845	aten.kthvalue.dimname,	845	aten.kthvalue.dimname,
846	aten.kthvalue.dimname_out,	846	aten.kthvalue.dimname_out,
847	aten.kthvalue.values,	847	aten.kthvalue.values,
848	aten.linalg_cholesky_ex.L,	848	aten.linalg_cholesky_ex.L,
849	aten.linalg_cholesky_ex.default,	849	aten.linalg_cholesky_ex.default,
850	aten.linalg_householder_product.default,	850	aten.linalg_householder_product.default,
851	aten.linalg_householder_product.out,	851	aten.linalg_householder_product.out,
852	aten.linalg_inv_ex.default,	852	aten.linalg_inv_ex.default,
853	aten.linalg_inv_ex.inverse,	853	aten.linalg_inv_ex.inverse,
854	aten.linalg_ldl_factor_ex.default,	854	aten.linalg_ldl_factor_ex.default,
855	aten.linalg_ldl_factor_ex.out,	855	aten.linalg_ldl_factor_ex.out,
856	aten.linalg_ldl_solve.default,	856	aten.linalg_ldl_solve.default,
857	aten.linalg_ldl_solve.out,	857	aten.linalg_ldl_solve.out,
858	aten.linalg_lu.default,	858	aten.linalg_lu.default,
859	aten.linalg_lu.out,	859	aten.linalg_lu.out,
860	aten.linalg_lu_factor_ex.default,	860	aten.linalg_lu_factor_ex.default,
861	aten.linalg_lu_factor_ex.out,	861	aten.linalg_lu_factor_ex.out,
862	aten.linalg_lu_solve.default,	862	aten.linalg_lu_solve.default,
863	aten.linalg_lu_solve.out,	863	aten.linalg_lu_solve.out,
864	aten.linalg_matrix_exp.default,	864	aten.linalg_matrix_exp.default,
865	aten.linalg_matrix_exp.out,	865	aten.linalg_matrix_exp.out,
866	aten.linalg_pinv.atol_rtol_tensor,	866	aten.linalg_pinv.atol_rtol_tensor,
867	aten.linalg_qr.default,	867	aten.linalg_qr.default,
868	aten.linalg_qr.out,	868	aten.linalg_qr.out,
869	aten.linalg_solve_triangular.default,	869	aten.linalg_solve_triangular.default,
870	aten.linalg_solve_triangular.out,	870	aten.linalg_solve_triangular.out,
871	aten.logcumsumexp.default,	871	aten.logcumsumexp.default,
872	aten.lu_unpack.default,	872	aten.lu_unpack.default,
873	aten.lu_unpack.out,	873	aten.lu_unpack.out,
874	aten.masked_scatter.default,	874	aten.masked_scatter.default,
875	aten.masked_scatter.out,	875	aten.masked_scatter.out,
876	aten.masked_scatter_backward.default,	876	aten.masked_scatter_backward.default,
877	aten.max_pool3d_with_indices.default,	877	aten.max_pool3d_with_indices.default,
878	aten.max_pool3d_with_indices.out,	878	aten.max_pool3d_with_indices.out,
879	aten.max_pool3d_with_indices_backward.default,	879	aten.max_pool3d_with_indices_backward.default,
880	aten.max_pool3d_with_indices_backward.grad_input,	880	aten.max_pool3d_with_indices_backward.grad_input,
881	aten.median.default,	881	aten.median.default,
882	aten.median.dim,	882	aten.median.dim,
883	aten.median.dim_values,	883	aten.median.dim_values,
884	aten.median.names_dim,	884	aten.median.names_dim,
885	aten.median.names_dim_values,	885	aten.median.names_dim_values,
886	aten.median.out,	886	aten.median.out,
887	aten.mode.default,	887	aten.mode.default,
888	aten.mode.dimname,	888	aten.mode.dimname,
889	aten.mode.dimname_out,	889	aten.mode.dimname_out,
890	aten.mode.values,	890	aten.mode.values,
891	aten.nanmedian.default,	891	aten.nanmedian.default,
892	aten.nanmedian.dim,	892	aten.nanmedian.dim,
893	aten.nanmedian.dim_values,	893	aten.nanmedian.dim_values,
894	aten.nanmedian.names_dim,	894	aten.nanmedian.names_dim,
895	aten.nanmedian.names_dim_values,	895	aten.nanmedian.names_dim_values,
896	aten.nanmedian.out,	896	aten.nanmedian.out,
897	aten.nonzero.default,	897	aten.nonzero.default,
898	aten.ormqr.default,	898	aten.ormqr.default,
899	aten.ormqr.out,	899	aten.ormqr.out,
900	aten.polygamma.default,	900	aten.polygamma.default,
901	aten.polygamma.out,	901	aten.polygamma.out,
902	aten.rand.default,	902	aten.rand.default,
903	aten.rand.generator,	903	aten.rand.generator,
904	aten.randint.default,	904	aten.randint.default,
905	aten.randint.generator,	905	aten.randint.generator,
906	aten.randint.generator_out,	906	aten.randint.generator_out,
907	aten.randint.low,	907	aten.randint.low,
908	aten.randint.low_generator,	908	aten.randint.low_generator,
909	aten.randint.low_generator_out,	909	aten.randint.low_generator_out,
910	aten.randint.low_out,	910	aten.randint.low_out,
911	aten.randint.out,	911	aten.randint.out,
912	aten.randn.default,	912	aten.randn.default,
913	aten.randn.generator,	913	aten.randn.generator,
914	aten.randperm.default,	914	aten.randperm.default,
915	aten.randperm.generator,	915	aten.randperm.generator,
916	aten.randperm.generator_out,	916	aten.randperm.generator_out,
917	aten.randperm.out,	917	aten.randperm.out,
918	aten.replication_pad1d_backward.default,	918	aten.replication_pad1d_backward.default,
919	aten.replication_pad1d_backward.grad_input,	919	aten.replication_pad1d_backward.grad_input,
920	aten.replication_pad2d_backward.default,	920	aten.replication_pad2d_backward.default,
921	aten.replication_pad2d_backward.grad_input,	921	aten.replication_pad2d_backward.grad_input,
922	aten.resize_.default,	922	aten.resize_.default,
923	aten.resize_as_.default,	923	aten.resize_as_.default,
924	aten.segment_reduce.default,	924	aten.segment_reduce.default,
925	aten.soft_margin_loss_backward.default,	925	aten.soft_margin_loss_backward.default,
926	aten.soft_margin_loss_backward.grad_input,	926	aten.soft_margin_loss_backward.grad_input,
927	aten.sort.Tensor,	927	aten.sort.Tensor,
928	aten.sort.any,	928	aten.sort.any,
929	aten.sort.bool,	929	aten.sort.bool,
930	aten.sort.default,	930	aten.sort.default,
931	aten.sort.dimname,	931	aten.sort.dimname,
932	aten.sort.dimname_stable,	932	aten.sort.dimname_stable,
933	aten.sort.dimname_values,	933	aten.sort.dimname_values,
934	aten.sort.dimname_values_stable,	934	aten.sort.dimname_values_stable,
935	aten.sort.float,	935	aten.sort.float,
936	aten.sort.int,	936	aten.sort.int,
937	aten.sort.stable,	937	aten.sort.stable,
938	aten.sort.str,	938	aten.sort.str,
939	aten.sort.values,	939	aten.sort.values,
940	aten.sort.values_stable,	940	aten.sort.values_stable,
941	aten.special_airy_ai.default,	941	aten.special_airy_ai.default,
942	aten.special_airy_ai.out,	942	aten.special_airy_ai.out,
943	aten.special_chebyshev_polynomial_t.default,	943	aten.special_chebyshev_polynomial_t.default,
944	aten.special_chebyshev_polynomial_t.n_scalar,	944	aten.special_chebyshev_polynomial_t.n_scalar,
945	aten.special_chebyshev_polynomial_t.n_scalar_out,	945	aten.special_chebyshev_polynomial_t.n_scalar_out,
946	aten.special_chebyshev_polynomial_t.out,	946	aten.special_chebyshev_polynomial_t.out,
947	aten.special_chebyshev_polynomial_t.x_scalar,	947	aten.special_chebyshev_polynomial_t.x_scalar,
948	aten.special_chebyshev_polynomial_t.x_scalar_out,	948	aten.special_chebyshev_polynomial_t.x_scalar_out,
949	aten.special_chebyshev_polynomial_u.default,	949	aten.special_chebyshev_polynomial_u.default,
950	aten.special_chebyshev_polynomial_u.n_scalar,	950	aten.special_chebyshev_polynomial_u.n_scalar,
951	aten.special_chebyshev_polynomial_u.n_scalar_out,	951	aten.special_chebyshev_polynomial_u.n_scalar_out,
952	aten.special_chebyshev_polynomial_u.out,	952	aten.special_chebyshev_polynomial_u.out,
953	aten.special_chebyshev_polynomial_u.x_scalar,	953	aten.special_chebyshev_polynomial_u.x_scalar,
954	aten.special_chebyshev_polynomial_u.x_scalar_out,	954	aten.special_chebyshev_polynomial_u.x_scalar_out,
955	aten.special_chebyshev_polynomial_v.default,	955	aten.special_chebyshev_polynomial_v.default,
956	aten.special_chebyshev_polynomial_v.n_scalar,	956	aten.special_chebyshev_polynomial_v.n_scalar,
957	aten.special_chebyshev_polynomial_v.n_scalar_out,	957	aten.special_chebyshev_polynomial_v.n_scalar_out,
958	aten.special_chebyshev_polynomial_v.out,	958	aten.special_chebyshev_polynomial_v.out,
959	aten.special_chebyshev_polynomial_v.x_scalar,	959	aten.special_chebyshev_polynomial_v.x_scalar,
960	aten.special_chebyshev_polynomial_v.x_scalar_out,	960	aten.special_chebyshev_polynomial_v.x_scalar_out,
961	aten.special_chebyshev_polynomial_w.default,	961	aten.special_chebyshev_polynomial_w.default,
962	aten.special_chebyshev_polynomial_w.n_scalar,	962	aten.special_chebyshev_polynomial_w.n_scalar,
963	aten.special_chebyshev_polynomial_w.n_scalar_out,	963	aten.special_chebyshev_polynomial_w.n_scalar_out,
964	aten.special_chebyshev_polynomial_w.out,	964	aten.special_chebyshev_polynomial_w.out,
965	aten.special_chebyshev_polynomial_w.x_scalar,	965	aten.special_chebyshev_polynomial_w.x_scalar,
966	aten.special_chebyshev_polynomial_w.x_scalar_out,	966	aten.special_chebyshev_polynomial_w.x_scalar_out,
967	aten.special_gammainc.default,	967	aten.special_gammainc.default,
968	aten.special_gammainc.out,	968	aten.special_gammainc.out,
969	aten.special_gammaincc.default,	969	aten.special_gammaincc.default,
970	aten.special_gammaincc.out,	970	aten.special_gammaincc.out,
971	aten.special_hermite_polynomial_h.default,	971	aten.special_hermite_polynomial_h.default,
972	aten.special_hermite_polynomial_h.n_scalar,	972	aten.special_hermite_polynomial_h.n_scalar,
973	aten.special_hermite_polynomial_h.n_scalar_out,	973	aten.special_hermite_polynomial_h.n_scalar_out,
974	aten.special_hermite_polynomial_h.out,	974	aten.special_hermite_polynomial_h.out,
975	aten.special_hermite_polynomial_h.x_scalar,	975	aten.special_hermite_polynomial_h.x_scalar,
976	aten.special_hermite_polynomial_h.x_scalar_out,	976	aten.special_hermite_polynomial_h.x_scalar_out,
977	aten.special_hermite_polynomial_he.default,	977	aten.special_hermite_polynomial_he.default,
978	aten.special_hermite_polynomial_he.n_scalar,	978	aten.special_hermite_polynomial_he.n_scalar,
979	aten.special_hermite_polynomial_he.n_scalar_out,	979	aten.special_hermite_polynomial_he.n_scalar_out,
980	aten.special_hermite_polynomial_he.out,	980	aten.special_hermite_polynomial_he.out,
981	aten.special_hermite_polynomial_he.x_scalar,	981	aten.special_hermite_polynomial_he.x_scalar,
982	aten.special_hermite_polynomial_he.x_scalar_out,	982	aten.special_hermite_polynomial_he.x_scalar_out,
983	aten.special_i0e.default,	983	aten.special_i0e.default,
984	aten.special_i0e.out,	984	aten.special_i0e.out,
985	aten.special_i1e.default,	985	aten.special_i1e.default,
986	aten.special_i1e.out,	986	aten.special_i1e.out,
987	aten.special_laguerre_polynomial_l.default,	987	aten.special_laguerre_polynomial_l.default,
988	aten.special_laguerre_polynomial_l.n_scalar,	988	aten.special_laguerre_polynomial_l.n_scalar,
989	aten.special_laguerre_polynomial_l.n_scalar_out,	989	aten.special_laguerre_polynomial_l.n_scalar_out,
990	aten.special_laguerre_polynomial_l.out,	990	aten.special_laguerre_polynomial_l.out,
991	aten.special_laguerre_polynomial_l.x_scalar,	991	aten.special_laguerre_polynomial_l.x_scalar,
992	aten.special_laguerre_polynomial_l.x_scalar_out,	992	aten.special_laguerre_polynomial_l.x_scalar_out,
993	aten.special_legendre_polynomial_p.default,	993	aten.special_legendre_polynomial_p.default,
994	aten.special_legendre_polynomial_p.n_scalar,	994	aten.special_legendre_polynomial_p.n_scalar,
995	aten.special_legendre_polynomial_p.n_scalar_out,	995	aten.special_legendre_polynomial_p.n_scalar_out,
996	aten.special_legendre_polynomial_p.out,	996	aten.special_legendre_polynomial_p.out,
997	aten.special_legendre_polynomial_p.x_scalar,	997	aten.special_legendre_polynomial_p.x_scalar,
998	aten.special_legendre_polynomial_p.x_scalar_out,	998	aten.special_legendre_polynomial_p.x_scalar_out,
999	aten.special_log_ndtr.default,	999	aten.special_log_ndtr.default,
1000	aten.special_log_ndtr.out,	1000	aten.special_log_ndtr.out,
1001	aten.special_modified_bessel_k0.default,	1001	aten.special_modified_bessel_k0.default,
1002	aten.special_modified_bessel_k0.out,	1002	aten.special_modified_bessel_k0.out,
1003	aten.special_modified_bessel_k1.default,	1003	aten.special_modified_bessel_k1.default,
1004	aten.special_modified_bessel_k1.out,	1004	aten.special_modified_bessel_k1.out,
1005	aten.special_ndtr.default,	1005	aten.special_ndtr.default,
1006	aten.special_ndtr.out,	1006	aten.special_ndtr.out,
1007	aten.special_ndtri.default,	1007	aten.special_ndtri.default,
1008	aten.special_ndtri.out,	1008	aten.special_ndtri.out,
1009	aten.special_scaled_modified_bessel_k0.default,	1009	aten.special_scaled_modified_bessel_k0.default,
1010	aten.special_scaled_modified_bessel_k0.out,	1010	aten.special_scaled_modified_bessel_k0.out,
1011	aten.special_scaled_modified_bessel_k1.default,	1011	aten.special_scaled_modified_bessel_k1.default,
1012	aten.special_scaled_modified_bessel_k1.out,	1012	aten.special_scaled_modified_bessel_k1.out,
1013	aten.special_shifted_chebyshev_polynomial_t.default,	1013	aten.special_shifted_chebyshev_polynomial_t.default,
1014	aten.special_shifted_chebyshev_polynomial_t.n_scalar,	1014	aten.special_shifted_chebyshev_polynomial_t.n_scalar,
1015	aten.special_shifted_chebyshev_polynomial_t.n_scalar_out,	1015	aten.special_shifted_chebyshev_polynomial_t.n_scalar_out,
1016	aten.special_shifted_chebyshev_polynomial_t.out,	1016	aten.special_shifted_chebyshev_polynomial_t.out,
1017	aten.special_shifted_chebyshev_polynomial_t.x_scalar,	1017	aten.special_shifted_chebyshev_polynomial_t.x_scalar,
1018	aten.special_shifted_chebyshev_polynomial_t.x_scalar_out,	1018	aten.special_shifted_chebyshev_polynomial_t.x_scalar_out,
1019	aten.special_shifted_chebyshev_polynomial_u.default,	1019	aten.special_shifted_chebyshev_polynomial_u.default,
1020	aten.special_shifted_chebyshev_polynomial_u.n_scalar,	1020	aten.special_shifted_chebyshev_polynomial_u.n_scalar,
1021	aten.special_shifted_chebyshev_polynomial_u.n_scalar_out,	1021	aten.special_shifted_chebyshev_polynomial_u.n_scalar_out,
1022	aten.special_shifted_chebyshev_polynomial_u.out,	1022	aten.special_shifted_chebyshev_polynomial_u.out,
1023	aten.special_shifted_chebyshev_polynomial_u.x_scalar,	1023	aten.special_shifted_chebyshev_polynomial_u.x_scalar,
1024	aten.special_shifted_chebyshev_polynomial_u.x_scalar_out,	1024	aten.special_shifted_chebyshev_polynomial_u.x_scalar_out,
1025	aten.special_shifted_chebyshev_polynomial_v.default,	1025	aten.special_shifted_chebyshev_polynomial_v.default,
1026	aten.special_shifted_chebyshev_polynomial_v.n_scalar,	1026	aten.special_shifted_chebyshev_polynomial_v.n_scalar,
1027	aten.special_shifted_chebyshev_polynomial_v.n_scalar_out,	1027	aten.special_shifted_chebyshev_polynomial_v.n_scalar_out,
1028	aten.special_shifted_chebyshev_polynomial_v.out,	1028	aten.special_shifted_chebyshev_polynomial_v.out,
1029	aten.special_shifted_chebyshev_polynomial_v.x_scalar,	1029	aten.special_shifted_chebyshev_polynomial_v.x_scalar,
1030	aten.special_shifted_chebyshev_polynomial_v.x_scalar_out,	1030	aten.special_shifted_chebyshev_polynomial_v.x_scalar_out,
1031	aten.special_shifted_chebyshev_polynomial_w.default,	1031	aten.special_shifted_chebyshev_polynomial_w.default,
1032	aten.special_shifted_chebyshev_polynomial_w.n_scalar,	1032	aten.special_shifted_chebyshev_polynomial_w.n_scalar,
1033	aten.special_shifted_chebyshev_polynomial_w.n_scalar_out,	1033	aten.special_shifted_chebyshev_polynomial_w.n_scalar_out,
1034	aten.special_shifted_chebyshev_polynomial_w.out,	1034	aten.special_shifted_chebyshev_polynomial_w.out,
1035	aten.special_shifted_chebyshev_polynomial_w.x_scalar,	1035	aten.special_shifted_chebyshev_polynomial_w.x_scalar,
1036	aten.special_shifted_chebyshev_polynomial_w.x_scalar_out,	1036	aten.special_shifted_chebyshev_polynomial_w.x_scalar_out,
1037	aten.special_spherical_bessel_j0.default,	1037	aten.special_spherical_bessel_j0.default,
1038	aten.special_spherical_bessel_j0.out,	1038	aten.special_spherical_bessel_j0.out,
1039	aten.special_zeta.default,	1039	aten.special_zeta.default,
1040	aten.special_zeta.other_scalar,	1040	aten.special_zeta.other_scalar,
1041	aten.special_zeta.other_scalar_out,	1041	aten.special_zeta.other_scalar_out,
1042	aten.special_zeta.out,	1042	aten.special_zeta.out,
1043	aten.special_zeta.self_scalar,	1043	aten.special_zeta.self_scalar,
1044	aten.special_zeta.self_scalar_out,	1044	aten.special_zeta.self_scalar_out,
1045	aten.to_sparse.default,	1045	aten.to_sparse.default,
1046	aten.to_sparse.sparse_dim,	1046	aten.to_sparse.sparse_dim,
1047	aten.topk.default,	1047	aten.topk.default,
1048	aten.topk.values,	1048	aten.topk.values,
1049	aten.triangular_solve.X,	1049	aten.triangular_solve.X,
1050	aten.triangular_solve.default,	1050	aten.triangular_solve.default,
1051	aten.uniform.default,	1051	aten.uniform.default,
1052	aten.uniform.out,	1052	aten.uniform.out,
1053	aten.upsample_bicubic2d_backward.default,	1053	aten.upsample_bicubic2d_backward.default,
1054	aten.upsample_bicubic2d_backward.grad_input,	1054	aten.upsample_bicubic2d_backward.grad_input,
1055	aten.upsample_linear1d_backward.default,	1055	aten.upsample_linear1d_backward.default,
1056	aten.upsample_linear1d_backward.grad_input,	1056	aten.upsample_linear1d_backward.grad_input,
1057	aten.upsample_trilinear3d_backward.default,	1057	aten.upsample_trilinear3d_backward.default,
1058	aten.upsample_trilinear3d_backward.grad_input,	1058	aten.upsample_trilinear3d_backward.grad_input,
1059	aten.view_as_complex.default,	1059	aten.view_as_complex.default,
1060	aten.zeros.names,	1060	aten.zeros.names,
1061	auto_functionalized,	1061	auto_functionalized,
1062	graphsafe_run_with_rng_state,	1062	graphsafe_run_with_rng_state,
1063	prims.digamma.default,	1063	prims.digamma.default,
1064	prims.igamma.default,	1064	prims.igamma.default,
1065	prims.igammac.default,	1065	prims.igammac.default,
1066	quantized.max_pool2d.default,	1066	quantized.max_pool2d.default,
1067	run_and_save_rng_state,	1067	run_and_save_rng_state,
1068	run_with_rng_state,	1068	run_with_rng_state,
1069	]	1069	]
1070		1070
1071	FALLBACK_LIST = TORCH_NATIVE_FALLBACK_LIST + NPU_EXTRA_FALLBACK_LIST	1071	FALLBACK_LIST = TORCH_NATIVE_FALLBACK_LIST + NPU_EXTRA_FALLBACK_LIST
1072		1072
1073	if inductor_indirect_memory_mode != 'linear':	1073	if inductor_indirect_memory_mode != 'linear':
1074	FALLBACK_LIST += [	1074	FALLBACK_LIST += [
1075	aten.isnan,	1075	aten.isnan,
1076	]	1076	]
1077		1077
1078	INDIRECT_MEM_FALLBACK_LIST = [	1078	INDIRECT_MEM_FALLBACK_LIST = [
1079	aten.cat,	1079	aten.cat,
1080	aten.embedding,	1080	aten.embedding,
1081	aten.embedding.default,	1081	aten.embedding.default,
1082	aten.embedding.out,	1082	aten.embedding.out,
1083	aten.gather,	1083	aten.gather,
1084	aten.gather.default,	1084	aten.gather.default,
1085	aten.gather.dimname_out,	1085	aten.gather.dimname_out,
1086	aten.gather.out,	1086	aten.gather.out,
1087	aten.index,	1087	aten.index,
1088	aten.index.Tensor,	1088	aten.index.Tensor,
1089	aten._unsafe_index,	1089	aten._unsafe_index,
1090	aten._unsafe_index.Tensor,	1090	aten._unsafe_index.Tensor,
1091	aten.index_put,	1091	aten.index_put,
1092	aten.index_put.default,	1092	aten.index_put.default,
1093	aten.index_put_,	1093	aten.index_put_,
1094	aten.index_put_.default,	1094	aten.index_put_.default,
1095	aten._unsafe_index_put,	1095	aten._unsafe_index_put,
1096	aten.scatter,	1096	aten.scatter,
1097	aten.scatter.src,	1097	aten.scatter.src,
1098	aten.scatter_,	1098	aten.scatter_,
1099	aten.scatter_.src,	1099	aten.scatter_.src,
1100	aten.scatter_reduce,	1100	aten.scatter_reduce,
1101	aten.scatter_reduce_,	1101	aten.scatter_reduce_,
1102	]	1102	]
1103		1103
1104	if not inductor_indirect_memory_mode:	1104	if not inductor_indirect_memory_mode:
1105	FALLBACK_LIST += INDIRECT_MEM_FALLBACK_LIST	1105	FALLBACK_LIST += INDIRECT_MEM_FALLBACK_LIST

Mtorch_npu/_inductor/tools/fallback_list_tool.py+16-16

	@@ -1,17 +1,17 @@
1	"""	1	"""
2	this tool is to get and save FALLBACK_LIST in lowering_fallback_list.py	2	this tool is to get and save FALLBACK_LIST in lowering_fallback_list.py
3	"""	3	"""
4		4
5	from torch_npu._inductor import lowering_fallback_list	5	from torch_npu._inductor import lowering_fallback_list
6	from torch_npu._inductor.tools.aten_op_tool import write_to_file	6	from torch_npu._inductor.tools.aten_op_tool import write_to_file
7		7
8		8
9	def get_npu_fallback_list():	9	def get_npu_fallback_list():
10	npu_fallback_list = {str(op) for op in lowering_fallback_list.FALLBACK_LIST}	10	npu_fallback_list = {str(op) for op in lowering_fallback_list.FALLBACK_LIST}
11	print(f"len(lowering_fallback_list.FALLBACK_LIST): {len(lowering_fallback_list.FALLBACK_LIST)}")	11	print(f"len(lowering_fallback_list.FALLBACK_LIST): {len(lowering_fallback_list.FALLBACK_LIST)}")
12	return npu_fallback_list	12	return npu_fallback_list
13		13
14		14
15	if __name__ == "__main__":	15	if __name__ == "__main__":
16	npu_fallback_list = get_npu_fallback_list()	16	npu_fallback_list = get_npu_fallback_list()
17	write_to_file(sorted(npu_fallback_list), f"npu_fallback_list_{len(npu_fallback_list)}.txt")	17	write_to_file(sorted(npu_fallback_list), f"npu_fallback_list_{len(npu_fallback_list)}.txt")

Mtorch_npu/csrc/afd/CMakeLists.txt+5-5

	@@ -1,6 +1,6 @@
1	FILE(GLOB _AFD_SRCS *.cpp)	1	FILE(GLOB _AFD_SRCS *.cpp)
2		2
3	LIST(APPEND AFD_SRCS ${_AFD_SRCS})	3	LIST(APPEND AFD_SRCS ${_AFD_SRCS})
4		4
5	# Pass to parent	5	# Pass to parent
6	set(AFD_SRCS ${AFD_SRCS} PARENT_SCOPE)	6	set(AFD_SRCS ${AFD_SRCS} PARENT_SCOPE)

Mtorch_npu/csrc/afd/Init.cpp+44-44

Mtorch_npu/csrc/afd/Init.h+11-11

	@@ -1,12 +1,12 @@
1	#pragma once	1	#pragma once
2		2
3	#include <torch/csrc/python_headers.h>	3	#include <torch/csrc/python_headers.h>
4	#include "torch_npu/csrc/core/npu/NPUMacros.h"	4	#include "torch_npu/csrc/core/npu/NPUMacros.h"
5		5
6	namespace torch_npu {	6	namespace torch_npu {
7	namespace afd {	7	namespace afd {
8		8
9	TORCH_NPU_API PyMethodDef *python_functions();	9	TORCH_NPU_API PyMethodDef *python_functions();
10		10
11	} // namespace afd	11	} // namespace afd
12	} // namespace torch_npu	12	} // namespace torch_npu

Mtorch_npu/csrc/aten/CMakeLists.txt+21-21

	@@ -1,21 +1,21 @@
1	FILE(GLOB _ATEN_SRCS	1	FILE(GLOB _ATEN_SRCS
2	*.cpp	2	*.cpp
3	common/*.cpp	3	common/*.cpp
4	mirror/*.cpp	4	mirror/*.cpp
5	ops/*.cpp	5	ops/*.cpp
6	ops/op_api/*.cpp)	6	ops/op_api/*.cpp)
7		7
8	FILE(GLOB _EXCLUDE	8	FILE(GLOB _EXCLUDE
9	VariableTypeEverything.cpp	9	VariableTypeEverything.cpp
10	ADInplaceOrViewTypeEverything.cpp	10	ADInplaceOrViewTypeEverything.cpp
11	python_functionsEverything.cpp	11	python_functionsEverything.cpp
12	RegisterFunctionalizationEverything.cpp)	12	RegisterFunctionalizationEverything.cpp)
13		13
14	FOREACH(ITEM ${_EXCLUDE})	14	FOREACH(ITEM ${_EXCLUDE})
15	LIST(REMOVE_ITEM _ATEN_SRCS ${ITEM})	15	LIST(REMOVE_ITEM _ATEN_SRCS ${ITEM})
16	ENDFOREACH()	16	ENDFOREACH()
17		17
18	LIST(APPEND ATEN_SRCS ${_ATEN_SRCS})	18	LIST(APPEND ATEN_SRCS ${_ATEN_SRCS})
19		19
20	# Pass to parent	20	# Pass to parent
21	set(ATEN_SRCS ${ATEN_SRCS} PARENT_SCOPE)	21	set(ATEN_SRCS ${ATEN_SRCS} PARENT_SCOPE)

Mtorch_npu/csrc/aten/PinMemory.cpp+50-50

	@@ -1,50 +1,50 @@
1	#include <ATen/ATen.h>	1	#include <ATen/ATen.h>
2	#include <ATen/NativeFunctions.h>	2	#include <ATen/NativeFunctions.h>
3	#include <ATen/TensorUtils.h>	3	#include <ATen/TensorUtils.h>
4	#include <c10/core/Storage.h>	4	#include <c10/core/Storage.h>
5	#include "torch_npu/csrc/core/npu/NPUFunctions.h"	5	#include "torch_npu/csrc/core/npu/NPUFunctions.h"
6	#include "torch_npu/csrc/core/npu/NPUException.h"	6	#include "torch_npu/csrc/core/npu/NPUException.h"
7	#include "torch_npu/csrc/core/npu/CachingHostAllocator.h"	7	#include "torch_npu/csrc/core/npu/CachingHostAllocator.h"
8		8
9	#define TORCH_ASSERT_ONLY_METHOD_OPERATORS	9	#define TORCH_ASSERT_ONLY_METHOD_OPERATORS
10	#include <ATen/core/Tensor.h>	10	#include <ATen/core/Tensor.h>
11	#include <ATen/core/dispatch/DispatchKeyExtractor.h>	11	#include <ATen/core/dispatch/DispatchKeyExtractor.h>
12	#include <torch/library.h>	12	#include <torch/library.h>
13		13
14	#ifndef AT_PER_OPERATOR_HEADERS	14	#ifndef AT_PER_OPERATOR_HEADERS
15	#include <ATen/Operators.h>	15	#include <ATen/Operators.h>
16	#else	16	#else
17	#include <ATen/ops/is_pinned_ops.h>	17	#include <ATen/ops/is_pinned_ops.h>
18	#include <ATen/ops/_pin_memory_ops.h>	18	#include <ATen/ops/_pin_memory_ops.h>
19		19
20	${ops_headers}	20	${ops_headers}
21	#endif	21	#endif
22		22
23	namespace at_npu {	23	namespace at_npu {
24	namespace native {	24	namespace native {
25		25
26	bool is_pinned(const at::Tensor& self, c10::optional<at::Device> device)	26	bool is_pinned(const at::Tensor& self, c10::optional<at::Device> device)
27	{	27	{
28	// Only CPU tensors can be pinned	28	// Only CPU tensors can be pinned
29	if (!self.is_cpu()) {	29	if (!self.is_cpu()) {
30	return false;	30	return false;
31	}	31	}
32	c10::DispatchKeySet _dk = c10::DispatchKeySet(c10::computeDispatchKey(c10::nullopt, self.layout(), device.value_or(c10::DeviceType::PrivateUse1)));	32	c10::DispatchKeySet _dk = c10::DispatchKeySet(c10::computeDispatchKey(c10::nullopt, self.layout(), device.value_or(c10::DeviceType::PrivateUse1)));
33	return at::_ops::is_pinned::redispatch(_dk, self, device);	33	return at::_ops::is_pinned::redispatch(_dk, self, device);
34	}	34	}
35		35
36	at::Tensor _pin_memory(const at::Tensor& self, c10::optional<at::Device> device)	36	at::Tensor _pin_memory(const at::Tensor& self, c10::optional<at::Device> device)
37	{	37	{
38	TORCH_CHECK(self.device().is_cpu(), "cannot pin '", self.toString(), "' only dense CPU tensors can be pinned", PTA_ERROR(ErrCode::TYPE));	38	TORCH_CHECK(self.device().is_cpu(), "cannot pin '", self.toString(), "' only dense CPU tensors can be pinned", PTA_ERROR(ErrCode::TYPE));
39	c10::DispatchKeySet _dk = c10::DispatchKeySet(c10::computeDispatchKey(c10::nullopt, self.layout(), device.value_or(c10::DeviceType::PrivateUse1)));	39	c10::DispatchKeySet _dk = c10::DispatchKeySet(c10::computeDispatchKey(c10::nullopt, self.layout(), device.value_or(c10::DeviceType::PrivateUse1)));
40	return at::_ops::_pin_memory::redispatch(_dk, self, device);	40	return at::_ops::_pin_memory::redispatch(_dk, self, device);
41	}	41	}
42		42
43	TORCH_LIBRARY_IMPL(aten, BackendSelect, m)	43	TORCH_LIBRARY_IMPL(aten, BackendSelect, m)
44	{	44	{
45	m.impl(TORCH_SELECTIVE_NAME("aten::is_pinned"), TORCH_FN(is_pinned));	45	m.impl(TORCH_SELECTIVE_NAME("aten::is_pinned"), TORCH_FN(is_pinned));
46	m.impl(TORCH_SELECTIVE_NAME("aten::_pin_memory"), TORCH_FN(_pin_memory));	46	m.impl(TORCH_SELECTIVE_NAME("aten::_pin_memory"), TORCH_FN(_pin_memory));
47	}	47	}
48		48
49	}	49	}
50	}	50	}

Mtorch_npu/csrc/aten/ops/op_api/CloneKernelOpApi.cpp+0-1

	@@ -45,4 +45,3 @@ at::Tensor NPUNativeOpApiFunctions::clone(const at::Tensor &src, c10::optional<c
45		45
46	} // namespace native	46	} // namespace native
47	} // namespace at_npu	47	} // namespace at_npu
48

Mtorch_npu/csrc/core/CMakeLists.txt+5-5

	@@ -1,6 +1,6 @@
1	FILE(GLOB _CORE_SRCS .cpp npu/.cpp npu//.cpp)	1	FILE(GLOB _CORE_SRCS .cpp npu/.cpp npu//.cpp)
2		2
3	LIST(APPEND CORE_SRCS ${_CORE_SRCS})	3	LIST(APPEND CORE_SRCS ${_CORE_SRCS})
4		4
5	# Pass to parent	5	# Pass to parent
6	set(CORE_SRCS ${CORE_SRCS} PARENT_SCOPE)	6	set(CORE_SRCS ${CORE_SRCS} PARENT_SCOPE)

Mtorch_npu/csrc/core/NPUStorageImpl.h+0-1

	@@ -67,4 +67,3 @@ c10::intrusive_ptr<c10::StorageImpl> make_npu_storage_impl(
67	bool resizable);	67	bool resizable);
68		68
69	} // namespace torch_npu	69	} // namespace torch_npu
70

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

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

	@@ -115,4 +115,3 @@ bool IsAclnnOnly()
115	return GetSocVersion() >= SocVersion::Ascend950;	115	return GetSocVersion() >= SocVersion::Ascend950;
116	}	116	}
117	} // namespace c10_npu	117	} // namespace c10_npu
118

Mtorch_npu/csrc/core/npu/NpuVariables.h+0-1

	@@ -45,4 +45,3 @@ bool IsBF16Supported();
45	bool IsAclnnOnly();	45	bool IsAclnnOnly();
46	} // namespace c10_npu	46	} // namespace c10_npu
47	#endif	47	#endif
48

Mtorch_npu/csrc/distributed/CMakeLists.txt+11-11

	@@ -1,12 +1,12 @@
1	if (DEFINED BUILD_LIBTORCH)	1	if (DEFINED BUILD_LIBTORCH)
2	FILE(GLOB _DIST_SRCS *.cpp)	2	FILE(GLOB _DIST_SRCS *.cpp)
3	# Exclude Python binding files when building libtorch	3	# Exclude Python binding files when building libtorch
4	list(REMOVE_ITEM _DIST_SRCS "${CMAKE_CURRENT_SOURCE_DIR}/Init.cpp")	4	list(REMOVE_ITEM _DIST_SRCS "${CMAKE_CURRENT_SOURCE_DIR}/Init.cpp")
5	else()	5	else()
6	FILE(GLOB _DIST_SRCS .cpp rpc/.cpp symm_mem/*.cpp)	6	FILE(GLOB _DIST_SRCS .cpp rpc/.cpp symm_mem/*.cpp)
7	endif()	7	endif()
8		8
9	LIST(APPEND DIST_SRCS ${_DIST_SRCS})	9	LIST(APPEND DIST_SRCS ${_DIST_SRCS})
10		10
11	# Pass to parent	11	# Pass to parent
12	set(DIST_SRCS ${DIST_SRCS} PARENT_SCOPE)	12	set(DIST_SRCS ${DIST_SRCS} PARENT_SCOPE)

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

Mtorch_npu/csrc/distributed/HcclCompile.h+496-496

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

Mtorch_npu/csrc/distributed/Init.h+11-11

	@@ -1,12 +1,12 @@
1	#pragma once	1	#pragma once
2		2
3	#include <torch/csrc/python_headers.h>	3	#include <torch/csrc/python_headers.h>
4	#include "torch_npu/csrc/core/npu/NPUMacros.h"	4	#include "torch_npu/csrc/core/npu/NPUMacros.h"
5		5
6	namespace torch_npu {	6	namespace torch_npu {
7	namespace distributed {	7	namespace distributed {
8		8
9	TORCH_NPU_API PyMethodDef* python_functions();	9	TORCH_NPU_API PyMethodDef* python_functions();
10		10
11	} // namespace distributed	11	} // namespace distributed
12	} // namespace torch_npu	12	} // namespace torch_npu

Mtorch_npu/csrc/distributed/default_comm_hooks.cpp+57-57

Mtorch_npu/csrc/distributed/default_comm_hooks.hpp+48-48

Mtorch_npu/csrc/distributed/reducer.cpp+2161-2161

Mtorch_npu/csrc/distributed/reducer.hpp+622-622

Mtorch_npu/csrc/framework/CMakeLists.txt+11-11

	@@ -1,12 +1,12 @@
1	FILE(GLOB _FRAMEWORK_SRCS	1	FILE(GLOB _FRAMEWORK_SRCS
2	*.cpp	2	*.cpp
3	aoe/*.cpp	3	aoe/*.cpp
4	autograd/*.cpp	4	autograd/*.cpp
5	contiguous/*.cpp	5	contiguous/*.cpp
6	interface/*.cpp	6	interface/*.cpp
7	utils/*.cpp)	7	utils/*.cpp)
8		8
9	LIST(APPEND FRAMEWORK_SRCS ${_FRAMEWORK_SRCS})	9	LIST(APPEND FRAMEWORK_SRCS ${_FRAMEWORK_SRCS})
10		10
11	# Pass to parent	11	# Pass to parent
12	set(FRAMEWORK_SRCS ${FRAMEWORK_SRCS} PARENT_SCOPE)	12	set(FRAMEWORK_SRCS ${FRAMEWORK_SRCS} PARENT_SCOPE)

Mtorch_npu/csrc/framework/StorageDescHelper.cpp+0-1

	@@ -311,4 +311,3 @@ int64_t StorageDescHelper::GetValidMemorySize(const at::Tensor &tensor)
311		311
312	} // namespace native	312	} // namespace native
313	} // namespace at_npu	313	} // namespace at_npu
314

Mtorch_npu/csrc/framework/contiguous/ReshapeOpt.cpp+0-1

	@@ -87,4 +87,3 @@ bool CanUseMemcpyForOtherFormat(const at::Tensor &tensor)
87		87
88	} // namespace native	88	} // namespace native
89	} // namespace at_npu	89	} // namespace at_npu
90

Mtorch_npu/csrc/framework/interface/LibAscendHal.h+0-1

	@@ -8,4 +8,3 @@ bool isSyscntEnable();
8	} // namespace native	8	} // namespace native
9	} // namespace torchat_npu_npu	9	} // namespace torchat_npu_npu
10	#endif	10	#endif
11

Mtorch_npu/csrc/npu/CMakeLists.txt+5-5

	@@ -1,6 +1,6 @@
1	FILE(GLOB _NPU_SRCS *.cpp)	1	FILE(GLOB _NPU_SRCS *.cpp)
2		2
3	LIST(APPEND NPU_SRCS ${_NPU_SRCS})	3	LIST(APPEND NPU_SRCS ${_NPU_SRCS})
4		4
5	# Pass to parent	5	# Pass to parent
6	set(NPU_SRCS ${NPU_SRCS} PARENT_SCOPE)	6	set(NPU_SRCS ${NPU_SRCS} PARENT_SCOPE)

Mtorch_npu/csrc/profiler/CMakeLists.txt+5-5

	@@ -1,6 +1,6 @@
1	FILE(GLOB _PROF_SRCS .cpp unwind/.cpp python/*.cpp)	1	FILE(GLOB _PROF_SRCS .cpp unwind/.cpp python/*.cpp)
2		2
3	LIST(APPEND PROF_SRCS ${_PROF_SRCS})	3	LIST(APPEND PROF_SRCS ${_PROF_SRCS})
4		4
5	# Pass to parent	5	# Pass to parent
6	set(PROF_SRCS ${PROF_SRCS} PARENT_SCOPE)	6	set(PROF_SRCS ${PROF_SRCS} PARENT_SCOPE)

Mtorch_npu/csrc/utils/TensorType.cpp+347-347

Mtorch_npu/csrc/utils/TensorType.h+16-16

	@@ -1,16 +1,16 @@
1	#include <torch/csrc/utils/tensor_new.h>	1	#include <torch/csrc/utils/tensor_new.h>
2		2
3	#include "torch_npu/csrc/core/npu/NPUMacros.h"	3	#include "torch_npu/csrc/core/npu/NPUMacros.h"
4	#include "torch_npu/csrc/core/npu/NPUFunctions.h"	4	#include "torch_npu/csrc/core/npu/NPUFunctions.h"
5		5
6	namespace torch_npu {	6	namespace torch_npu {
7	namespace utils {	7	namespace utils {
8		8
9	// Initializes the Python tensor type objects: torch.npu.FloatTensor,	9	// Initializes the Python tensor type objects: torch.npu.FloatTensor,
10	// torch.npu.DoubleTensor, etc. and binds them in their containing modules.	10	// torch.npu.DoubleTensor, etc. and binds them in their containing modules.
11	void _initialize_python_bindings();	11	void _initialize_python_bindings();
12		12
13	TORCH_NPU_API PyMethodDef* npu_extension_functions();	13	TORCH_NPU_API PyMethodDef* npu_extension_functions();
14		14
15	}	15	}
16	}	16	}

Mtorch_npu/distributed/__init__.py+30-30

Mtorch_npu/dynamo/trace_rule.py+99-100

Mtorch_npu/npu/amp/__init__.py+5-5

	@@ -1,6 +1,6 @@
1	__all__ = [	1	__all__ = [
2	"autocast", "GradScaler", "custom_fwd", "custom_bwd"	2	"autocast", "GradScaler", "custom_fwd", "custom_bwd"
3	]	3	]
4		4
5	from .autocast_mode import autocast, custom_fwd, custom_bwd # noqa: F401	5	from .autocast_mode import autocast, custom_fwd, custom_bwd # noqa: F401
		O O openLiBingCI5月17日此条代码评论区间+2至+5 【openlibing.ci】识别到代码检查告警抑制注释，匹配工具：flake8，请Committer检视其合理性。
6	from .grad_scaler import GradScaler # noqa: F401	6	from .grad_scaler import GradScaler # noqa: F401

Mtorch_npu/npu/amp/autocast_mode.py+97-97

	@@ -1,98 +1,98 @@
1	__all__ = ["autocast", "custom_fwd", "custom_bwd"]	1	__all__ = ["autocast", "custom_fwd", "custom_bwd"]
2		2
3		3
4	import functools	4	import functools
5	import collections	5	import collections
6	from typing import Any	6	from typing import Any
7	from typing_extensions import deprecated	7	from typing_extensions import deprecated
8		8
9	try:	9	try:
10	import numpy as np	10	import numpy as np
11		11
12	HAS_NUMPY = True	12	HAS_NUMPY = True
13	except ModuleNotFoundError:	13	except ModuleNotFoundError:
14	np = None # type: ignore[assignment]	14	np = None # type: ignore[assignment]
15		15
16	import torch	16	import torch
17	import torch_npu	17	import torch_npu
		O O openLiBingCI5月17日此条代码评论区间+12至+17 【openlibing.ci】识别到代码检查告警抑制注释，匹配工具：mypy，请Committer检视其合理性。
18	from torch_npu.utils._error_code import ErrCode, pta_error	18	from torch_npu.utils._error_code import ErrCode, pta_error
19		19
20		20
21	class autocast(torch.amp.autocast_mode.autocast):	21	class autocast(torch.amp.autocast_mode.autocast):
22	r"""	22	r"""
23	See :class:`torch.autocast`.	23	See :class:`torch.autocast`.
24	``torch.npu.amp.autocast(args...)`` is equivalent to ``torch.autocast("npu", args...)``	24	``torch.npu.amp.autocast(args...)`` is equivalent to ``torch.autocast("npu", args...)``
25	"""	25	"""
26		26
27	def __init__(self, enabled: bool = True, dtype: torch.dtype = torch.float16, cache_enabled: bool = True):	27	def __init__(self, enabled: bool = True, dtype: torch.dtype = torch.float16, cache_enabled: bool = True):
28	if torch._jit_internal.is_scripting():	28	if torch._jit_internal.is_scripting():
29	self._enabled = enabled	29	self._enabled = enabled
30	self.device = "npu"	30	self.device = "npu"
31	self.fast_dtype = dtype	31	self.fast_dtype = dtype
32	return	32	return
33	super().__init__("npu", enabled=enabled, dtype=dtype, cache_enabled=cache_enabled)	33	super().__init__("npu", enabled=enabled, dtype=dtype, cache_enabled=cache_enabled)
34		34
35	def __enter__(self):	35	def __enter__(self):
36	if torch._jit_internal.is_scripting():	36	if torch._jit_internal.is_scripting():
37	return self	37	return self
38	return super().__enter__()	38	return super().__enter__()
39		39
40	def __exit__(self, exc_type: Any, exc_val: Any, exc_tb: Any): # type: ignore[override]	40	def __exit__(self, exc_type: Any, exc_val: Any, exc_tb: Any): # type: ignore[override]
41	if torch._jit_internal.is_scripting():	41	if torch._jit_internal.is_scripting():
42	return None	42	return None
		O O openLiBingCI5月17日此条代码评论区间+37至+42 【openlibing.ci】识别到代码检查告警抑制注释，匹配工具：mypy，请Committer检视其合理性。
43	return super().__exit__(exc_type, exc_val, exc_tb)	43	return super().__exit__(exc_type, exc_val, exc_tb)
44		44
45	def __call__(self, func):	45	def __call__(self, func):
46	if torch._jit_internal.is_scripting():	46	if torch._jit_internal.is_scripting():
47	return func	47	return func
48	return super().__call__(func)	48	return super().__call__(func)
49		49
50		50
51	# Casts Tensors and containers of Tensors. Special-cases passthroughs for strings and np.ndarrays, which	51	# Casts Tensors and containers of Tensors. Special-cases passthroughs for strings and np.ndarrays, which
52	# may be falsely detected as "Iterables."	52	# may be falsely detected as "Iterables."
53	def _cast(value, dtype):	53	def _cast(value, dtype):
54	if isinstance(value, torch.Tensor):	54	if isinstance(value, torch.Tensor):
55	is_eligible = (value.is_floating_point() and value.device.type == 'npu' and (value.dtype is not torch.float64))	55	is_eligible = (value.is_floating_point() and value.device.type == 'npu' and (value.dtype is not torch.float64))
56	return value.to(dtype) if is_eligible else value	56	return value.to(dtype) if is_eligible else value
57	elif isinstance(value, (str, bytes)):	57	elif isinstance(value, (str, bytes)):
58	return value	58	return value
59	elif HAS_NUMPY and isinstance(value, np.ndarray):	59	elif HAS_NUMPY and isinstance(value, np.ndarray):
60	return value	60	return value
61	elif isinstance(value, collections.abc.Mapping):	61	elif isinstance(value, collections.abc.Mapping):
62	return {_cast(k, dtype): _cast(v, dtype) for k, v in value.items()}	62	return {_cast(k, dtype): _cast(v, dtype) for k, v in value.items()}
63	elif isinstance(value, collections.abc.Iterable):	63	elif isinstance(value, collections.abc.Iterable):
64	iterable = map(lambda v: _cast(v, dtype), value)	64	iterable = map(lambda v: _cast(v, dtype), value)
65	if isinstance(value, list) or isinstance(value, tuple):	65	if isinstance(value, list) or isinstance(value, tuple):
66	return type(value)(iterable)	66	return type(value)(iterable)
67	else:	67	else:
68	return iterable	68	return iterable
69	else:	69	else:
70	return value	70	return value
71		71
72		72
73	@deprecated(	73	@deprecated(
74	"`torch_npu.npu.amp.custom_fwd(args...)` is deprecated. "	74	"`torch_npu.npu.amp.custom_fwd(args...)` is deprecated. "
75	"Please use `torch.amp.custom_fwd(args..., device_type='npu')` instead.",	75	"Please use `torch.amp.custom_fwd(args..., device_type='npu')` instead.",
76	category=FutureWarning,	76	category=FutureWarning,
77	)	77	)
78	def custom_fwd(fwd=None, *, cast_inputs=None):	78	def custom_fwd(fwd=None, *, cast_inputs=None):
79	"""	79	"""
80	``torch_npu.npu.amp.custom_fwd(args...)`` is deprecated. Please use	80	``torch_npu.npu.amp.custom_fwd(args...)`` is deprecated. Please use
81	``torch.amp.custom_fwd(args..., device_type='npu')`` instead.	81	``torch.amp.custom_fwd(args..., device_type='npu')`` instead.
82	"""	82	"""
83	return functools.partial(torch.amp.custom_fwd, device_type="npu")(	83	return functools.partial(torch.amp.custom_fwd, device_type="npu")(
84	fwd=fwd, cast_inputs=cast_inputs	84	fwd=fwd, cast_inputs=cast_inputs
85	)	85	)
86		86
87		87
88	@deprecated(	88	@deprecated(
89	"`torch_npu.npu.amp.custom_bwd(args...)` is deprecated. "	89	"`torch_npu.npu.amp.custom_bwd(args...)` is deprecated. "
90	"Please use `torch.amp.custom_bwd(args..., device_type='npu')` instead.",	90	"Please use `torch.amp.custom_bwd(args..., device_type='npu')` instead.",
91	category=FutureWarning,	91	category=FutureWarning,
92	)	92	)
93	def custom_bwd(bwd):	93	def custom_bwd(bwd):
94	"""	94	"""
95	``torch_npu.npu.amp.custom_bwd(args...)`` is deprecated. Please use	95	``torch_npu.npu.amp.custom_bwd(args...)`` is deprecated. Please use
96	``torch.amp.custom_bwd(args..., device_type='npu')`` instead.	96	``torch.amp.custom_bwd(args..., device_type='npu')`` instead.
97	"""	97	"""
98	return functools.partial(torch.amp.custom_bwd, device_type="npu")(bwd)	98	return functools.partial(torch.amp.custom_bwd, device_type="npu")(bwd)

Mtorch_npu/npu/amp/common.py+5-5

	@@ -1,5 +1,5 @@
1	import torch_npu	1	import torch_npu
2		2
3		3
4	def amp_definitely_not_available():	4	def amp_definitely_not_available():
5	return not torch_npu.npu.is_available()	5	return not torch_npu.npu.is_available()

Mtorch_npu/npu/amp/grad_scaler.py+505-505

Mtorch_npu/testing/common_distributed.py+110-110

	@@ -1,110 +1,110 @@
1	import unittest	1	import unittest
2	from functools import wraps	2	from functools import wraps
3	from typing import (	3	from typing import (
4	Tuple,	4	Tuple,
5	Dict,	5	Dict,
6	Any,	6	Any,
7	)	7	)
8	from collections import namedtuple	8	from collections import namedtuple
9	import sys	9	import sys
10	import os	10	import os
11	from contextlib import contextmanager	11	from contextlib import contextmanager
12	import torch	12	import torch
13	import torch.distributed as dist	13	import torch.distributed as dist
14	import torch_npu	14	import torch_npu
15		15
16		16
17	TestSkip = namedtuple('TestSkip', 'exit_code, message')	17	TestSkip = namedtuple('TestSkip', 'exit_code, message')
18	TEST_SKIPS = {	18	TEST_SKIPS = {
19	"multi-npu": TestSkip(75, "Multi-NPU condition not satisfied"),	19	"multi-npu": TestSkip(75, "Multi-NPU condition not satisfied"),
20	"multi-npu-1": TestSkip(75, "Need at least 1 ASCEND devices"),	20	"multi-npu-1": TestSkip(75, "Need at least 1 ASCEND devices"),
21	"multi-npu-2": TestSkip(75, "Need at least 2 ASCEND devices"),	21	"multi-npu-2": TestSkip(75, "Need at least 2 ASCEND devices"),
22	"multi-npu-3": TestSkip(75, "Need at least 3 ASCEND devices"),	22	"multi-npu-3": TestSkip(75, "Need at least 3 ASCEND devices"),
23	"multi-npu-4": TestSkip(75, "Need at least 4 ASCEND devices"),	23	"multi-npu-4": TestSkip(75, "Need at least 4 ASCEND devices"),
24	"multi-npu-5": TestSkip(75, "Need at least 5 ASCEND devices"),	24	"multi-npu-5": TestSkip(75, "Need at least 5 ASCEND devices"),
25	"multi-npu-6": TestSkip(75, "Need at least 6 ASCEND devices"),	25	"multi-npu-6": TestSkip(75, "Need at least 6 ASCEND devices"),
26	"multi-npu-7": TestSkip(75, "Need at least 7 ASCEND devices"),	26	"multi-npu-7": TestSkip(75, "Need at least 7 ASCEND devices"),
27	"multi-npu-8": TestSkip(75, "Need at least 8 ASCEND devices"),	27	"multi-npu-8": TestSkip(75, "Need at least 8 ASCEND devices"),
28	"hccl":TestSkip(76, "c10d not compiled with HCCL support"),	28	"hccl":TestSkip(76, "c10d not compiled with HCCL support"),
29	"known_issues":TestSkip(77, "Test skipped due to known issues"),	29	"known_issues":TestSkip(77, "Test skipped due to known issues"),
30	}	30	}
31		31
32		32
33	def skipIfUnsupportMultiNPU(npu_number_needed):	33	def skipIfUnsupportMultiNPU(npu_number_needed):
34	def skip_dec(func):	34	def skip_dec(func):
35	@wraps(func)	35	@wraps(func)
36	def wrapper(self, args, *kwargs):	36	def wrapper(self, args, *kwargs):
37	if not torch.npu.is_available() or torch.npu.device_count() < npu_number_needed:	37	if not torch.npu.is_available() or torch.npu.device_count() < npu_number_needed:
38	raise unittest.SkipTest(f"Multi-NPU {npu_number_needed} condition not satisfied")	38	raise unittest.SkipTest(f"Multi-NPU {npu_number_needed} condition not satisfied")
39	return func(self, args, *kwargs)	39	return func(self, args, *kwargs)
40	return wrapper	40	return wrapper
41	return skip_dec	41	return skip_dec
42		42
43		43
44	def with_comms(func):	44	def with_comms(func):
45	if func is None:	45	if func is None:
46	raise RuntimeError("Test function is None.")	46	raise RuntimeError("Test function is None.")
47		47
48	@wraps(func) # pyre-ignore[6]	48	@wraps(func) # pyre-ignore[6]
49	def wrapper(	49	def wrapper(
50	self, args: Tuple[object], *kwargs: Dict[str, Any] # type: ignore[misc]	50	self, args: Tuple[object], *kwargs: Dict[str, Any] # type: ignore[misc]
51	) -> None:	51	) -> None:
52	# if backend not specified, and npu available, then use hccl, else gloo	52	# if backend not specified, and npu available, then use hccl, else gloo
		O O openLiBingCI5月17日此条代码评论区间+48至+52 【openlibing.ci】识别到代码检查告警抑制注释，匹配工具：mypy，请Committer检视其合理性。
53	if torch.npu.is_available() and torch.npu.device_count() >= self.world_size:	53	if torch.npu.is_available() and torch.npu.device_count() >= self.world_size:
54	self.device_type = "npu"	54	self.device_type = "npu"
55	else:	55	else:
56	self.device_type = "cpu"	56	self.device_type = "cpu"
57		57
58	pg_backend = (	58	pg_backend = (
59	"hccl" if self.device_type == "npu" else "gloo"	59	"hccl" if self.device_type == "npu" else "gloo"
60	)	60	)
61	if pg_backend == "hccl" and torch.npu.device_count() < self.world_size:	61	if pg_backend == "hccl" and torch.npu.device_count() < self.world_size:
62	raise RuntimeError(TEST_SKIPS[f"multi-npu-{self.world_size}"].message)	62	raise RuntimeError(TEST_SKIPS[f"multi-npu-{self.world_size}"].message)
63		63
64	init_pg(backend=pg_backend, world_size=self.world_size, rank=self.rank, file_name=self.file_name)	64	init_pg(backend=pg_backend, world_size=self.world_size, rank=self.rank, file_name=self.file_name)
65		65
66	torch.npu.manual_seed(0)	66	torch.npu.manual_seed(0)
67	torch.npu.initial_seed()	67	torch.npu.initial_seed()
68	func(self, args, *kwargs) # type: ignore[misc]	68	func(self, args, *kwargs) # type: ignore[misc]
69	self.destroy_pg()	69	self.destroy_pg()
70		70
71	return wrapper	71	return wrapper
		O O openLiBingCI5月17日此条代码评论区间+66至+71 【openlibing.ci】识别到代码检查告警抑制注释，匹配工具：mypy，请Committer检视其合理性。
72		72
73		73
74	def init_pg(backend: str = "hccl", world_size=1, rank=0, file_name="file://") -> None:	74	def init_pg(backend: str = "hccl", world_size=1, rank=0, file_name="file://") -> None:
75	if backend == "hccl" and torch.npu.device_count() < world_size:	75	if backend == "hccl" and torch.npu.device_count() < world_size:
76	raise RuntimeError(TEST_SKIPS[f"multi-npu-{world_size}"].message)	76	raise RuntimeError(TEST_SKIPS[f"multi-npu-{world_size}"].message)
77		77
78	if backend not in ["hccl", "gloo"]:	78	if backend not in ["hccl", "gloo"]:
79	raise RuntimeError(f"Backend {backend} not supported!")	79	raise RuntimeError(f"Backend {backend} not supported!")
80		80
81	dist.init_process_group(	81	dist.init_process_group(
82	backend=backend,	82	backend=backend,
83	world_size=world_size,	83	world_size=world_size,
84	rank=rank, # pyre-ignore[16]	84	rank=rank, # pyre-ignore[16]
85	init_method=f"file://{file_name}", # pyre-ignore[16]	85	init_method=f"file://{file_name}", # pyre-ignore[16]
86	)	86	)
87		87
88	# set device for hccl pg for collectives	88	# set device for hccl pg for collectives
89	if backend == "hccl":	89	if backend == "hccl":
90	torch.npu.set_device(rank)	90	torch.npu.set_device(rank)
91		91
92		92
93	@contextmanager	93	@contextmanager
94	def _dynamo_dist_per_rank_init(rank, world_size, init_pg_=True):	94	def _dynamo_dist_per_rank_init(rank, world_size, init_pg_=True):
95	# To avoid multiple inheritance from _dynamo.test_case.TestCase and MultiProcessTestCase,	95	# To avoid multiple inheritance from _dynamo.test_case.TestCase and MultiProcessTestCase,
96	# Just manually implement the most important part of the dynamo behavior to reset/clear.	96	# Just manually implement the most important part of the dynamo behavior to reset/clear.
97	torch_npu.npu.set_device(rank)	97	torch_npu.npu.set_device(rank)
98	os.environ['MASTER_ADDR'] = 'localhost'	98	os.environ['MASTER_ADDR'] = 'localhost'
99	os.environ['MASTER_PORT'] = '6789'	99	os.environ['MASTER_PORT'] = '6789'
100	if init_pg_:	100	if init_pg_:
101	dist.init_process_group(backend="hccl", rank=rank, world_size=world_size)	101	dist.init_process_group(backend="hccl", rank=rank, world_size=world_size)
102	torch._dynamo.reset()	102	torch._dynamo.reset()
103	torch._dynamo.utils.counters.clear()	103	torch._dynamo.utils.counters.clear()
104	try:	104	try:
105	yield	105	yield
106	finally:	106	finally:
107	torch._dynamo.reset()	107	torch._dynamo.reset()
108	torch._dynamo.utils.counters.clear()	108	torch._dynamo.utils.counters.clear()
109	if init_pg_:	109	if init_pg_:
110	dist.destroy_process_group()	110	dist.destroy_process_group()

Mtorch_npu/testing/decorator.py+142-142

	@@ -1,142 +1,142 @@
1	from functools import wraps, partialmethod	1	from functools import wraps, partialmethod
2		2
3	import os	3	import os
4	import inspect	4	import inspect
5	import itertools	5	import itertools
6	import torch	6	import torch
7		7
8		8
9	def feed_data(func, new_name, args, *kwargs):	9	def feed_data(func, new_name, args, *kwargs):
10	"""	10	"""
11	This internal method decorator feeds the test data item to the test.	11	This internal method decorator feeds the test data item to the test.
12	"""	12	"""
13	@wraps(func)	13	@wraps(func)
14	def wrapper(self):	14	def wrapper(self):
15	return func(self, args, *kwargs)	15	return func(self, args, *kwargs)
16	wrapper.__name__ = new_name	16	wrapper.__name__ = new_name
17	wrapper.__wrapped__ = func	17	wrapper.__wrapped__ = func
18	return wrapper	18	return wrapper
19		19
20		20
21	def instantiate_tests(arg=None, **kwargs):	21	def instantiate_tests(arg=None, **kwargs):
22		22
23	def wrapper(cls):	23	def wrapper(cls):
24	def gen_testcase(cls, func, name, key_list, func_args, value):	24	def gen_testcase(cls, func, name, key_list, func_args, value):
25	new_kwargs = dict(device="npu") if "device" in func_args else {}	25	new_kwargs = dict(device="npu") if "device" in func_args else {}
26	test_name = name	26	test_name = name
27	for k, v in zip(key_list, value):	27	for k, v in zip(key_list, value):
28	func_key = None	28	func_key = None
29	if k == "format":	29	if k == "format":
30	test_name += ("_" + str(v))	30	test_name += ("_" + str(v))
31	elif k == "dtype":	31	elif k == "dtype":
32	test_name += ("_" + str(v).split('.')[1])	32	test_name += ("_" + str(v).split('.')[1])
33	for _func_key in func_args:	33	for _func_key in func_args:
34	if k in _func_key:	34	if k in _func_key:
35	if func_key is not None:	35	if func_key is not None:
36	raise RuntimeError(f"Multiple matches for {k}")	36	raise RuntimeError(f"Multiple matches for {k}")
37	func_key = _func_key	37	func_key = _func_key
38	new_kwargs[func_key] = v	38	new_kwargs[func_key] = v
39	setattr(cls, test_name, feed_data(func, test_name, **new_kwargs))	39	setattr(cls, test_name, feed_data(func, test_name, **new_kwargs))
40		40
41	for name, func in list(cls.__dict__.items()):	41	for name, func in list(cls.__dict__.items()):
42	data = {}	42	data = {}
43	if hasattr(func, "dtypes"):	43	if hasattr(func, "dtypes"):
44	data['dtype'] = func.dtypes	44	data['dtype'] = func.dtypes
45	if hasattr(func, "formats"):	45	if hasattr(func, "formats"):
46	data['format'] = func.formats	46	data['format'] = func.formats
47		47
48	key_list = data.keys()	48	key_list = data.keys()
49	if not key_list:	49	if not key_list:
50	continue	50	continue
51		51
52	func_args = inspect.getfullargspec(func).args	52	func_args = inspect.getfullargspec(func).args
53	value_list = [data.get(key) for key in key_list]	53	value_list = [data.get(key) for key in key_list]
54	for value in itertools.product(*value_list):	54	for value in itertools.product(*value_list):
55	gen_testcase(cls, func, name, key_list, func_args, value)	55	gen_testcase(cls, func, name, key_list, func_args, value)
56		56
57	delattr(cls, name)	57	delattr(cls, name)
58	return cls	58	return cls
59		59
60	return wrapper(arg)	60	return wrapper(arg)
61		61
62		62
63	def gen_ops_testcase(cls, func, name, keys, value, op_info):	63	def gen_ops_testcase(cls, func, name, keys, value, op_info):
64	new_kwargs = {}	64	new_kwargs = {}
65	test_name = f'{func.__name__}_{name}'	65	test_name = f'{func.__name__}_{name}'
66		66
67	for k, v in zip(keys, value):	67	for k, v in zip(keys, value):
68	if k == "npu_format":	68	if k == "npu_format":
69	test_name += ("_" + str(v))	69	test_name += ("_" + str(v))
70	elif k == "dtype":	70	elif k == "dtype":
71	test_name += ("_" + str(v).split('.')[1])	71	test_name += ("_" + str(v).split('.')[1])
72	new_kwargs[k] = v	72	new_kwargs[k] = v
73		73
74	new_kwargs['op'] = op_info	74	new_kwargs['op'] = op_info
75	new_func = partialmethod(func, **new_kwargs)	75	new_func = partialmethod(func, **new_kwargs)
76		76
77	setattr(cls, test_name, new_func)	77	setattr(cls, test_name, new_func)
78	for decorator in op_info.get_decorators(cls.__name__, func.__name__, 'cpu', value[0], {}):	78	for decorator in op_info.get_decorators(cls.__name__, func.__name__, 'cpu', value[0], {}):
79	setattr(cls, test_name, decorator(new_func))	79	setattr(cls, test_name, decorator(new_func))
80		80
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
88	if 'test_variant_consistency_eager' in testcase:	88	if 'test_variant_consistency_eager' in testcase:
89	if torch.float32 in op_info.dtypesIfNPU:	89	if torch.float32 in op_info.dtypesIfNPU:
90	data['dtype'] = {torch.float32}	90	data['dtype'] = {torch.float32}
91	else:	91	else:
92	data['dtype'] = {list(op_info.dtypesIfNPU)[-1]}	92	data['dtype'] = {list(op_info.dtypesIfNPU)[-1]}
93		93
94	return data	94	return data
95		95
96		96
97	def instantiate_ops_tests(op_db):	97	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:
105	data = gen_op_input(testcase, func, op_info)	105	data = gen_op_input(testcase, func, op_info)
106	keys = data.keys()	106	keys = data.keys()
107	values = [data.get(key) for key in keys]	107	values = [data.get(key) for key in keys]
108		108
109	for value in itertools.product(*values):	109	for value in itertools.product(*values):
110	gen_ops_testcase(cls, func, op_info.name, keys, value, op_info)	110	gen_ops_testcase(cls, func, op_info.name, keys, value, op_info)
111		111
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
119	class Dtypes(object):	119	class Dtypes(object):
120		120
121	def __init__(self, *args):	121	def __init__(self, *args):
122	if (args is None or len(args) == 0):	122	if (args is None or len(args) == 0):
123	raise RuntimeError("No dtypes given")	123	raise RuntimeError("No dtypes given")
124	if not all(isinstance(arg, torch.dtype) for arg in args):	124	if not all(isinstance(arg, torch.dtype) for arg in args):
125	raise RuntimeError("Unknown dtype in {0}".format(str(args)))	125	raise RuntimeError("Unknown dtype in {0}".format(str(args)))
126	self.args = args	126	self.args = args
127		127
128	def __call__(self, fn):	128	def __call__(self, fn):
129	fn.dtypes = self.args	129	fn.dtypes = self.args
130	return fn	130	return fn
131		131
132		132
133	class Formats(object):	133	class Formats(object):
134		134
135	def __init__(self, *args):	135	def __init__(self, *args):
136	if args is None or len(args) == 0:	136	if args is None or len(args) == 0:
137	raise RuntimeError("No formats given")	137	raise RuntimeError("No formats given")
138	self.args = args	138	self.args = args
139		139
140	def __call__(self, fn):	140	def __call__(self, fn):
141	fn.formats = self.args	141	fn.formats = self.args
142	return fn	142	return fn

Mtorch_npu/utils/cpp_extension.py+60-60

	@@ -1,60 +1,60 @@
1	import os	1	import os
2	import setuptools	2	import setuptools
3		3
4	import torch	4	import torch
5	import torch.utils.cpp_extension as TorchExtension	5	import torch.utils.cpp_extension as TorchExtension
6		6
7	import torch_npu	7	import torch_npu
8		8
9	PYTORCH_NPU_INSTALL_PATH = os.path.dirname(os.path.realpath(torch_npu.__file__))	9	PYTORCH_NPU_INSTALL_PATH = os.path.dirname(os.path.realpath(torch_npu.__file__))
10		10
11		11
12	def NpuExtension(name, sources, args, *kwargs):	12	def NpuExtension(name, sources, args, *kwargs):
13	r'''	13	r'''
14	Creates a :class:`setuptools.Extension` for C++.	14	Creates a :class:`setuptools.Extension` for C++.
15		15
16	Convenience method that creates a :class:`setuptools.Extension` with the	16	Convenience method that creates a :class:`setuptools.Extension` with the
17	bare minimum (but often sufficient) arguments to build a C++ extension.	17	bare minimum (but often sufficient) arguments to build a C++ extension.
18		18
19	All arguments are forwarded to the :class:`setuptools.Extension`	19	All arguments are forwarded to the :class:`setuptools.Extension`
20	constructor.	20	constructor.
21		21
22	Example:	22	Example:
23	>>> from setuptools import setup	23	>>> from setuptools import setup
24	>>> from torch_npu.utils.cpp_extension import NpuExtension	24	>>> from torch_npu.utils.cpp_extension import NpuExtension
25	>>> setup(	25	>>> setup(
26	name='extension',	26	name='extension',
27	ext_modules=[	27	ext_modules=[
28	NpuExtension(	28	NpuExtension(
29	name='extension',	29	name='extension',
30	sources=['extension.cpp'],	30	sources=['extension.cpp'],
31	extra_compile_args=['-g']),	31	extra_compile_args=['-g']),
32	],	32	],
33	cmdclass={	33	cmdclass={
34	'build_ext': BuildExtension	34	'build_ext': BuildExtension
35	})	35	})
36	'''	36	'''
37		37
38	torch_npu_dir = PYTORCH_NPU_INSTALL_PATH	38	torch_npu_dir = PYTORCH_NPU_INSTALL_PATH
39	include_dirs = kwargs.get('include_dirs', [])	39	include_dirs = kwargs.get('include_dirs', [])
40	include_dirs.append(os.path.join(torch_npu_dir, 'include'))	40	include_dirs.append(os.path.join(torch_npu_dir, 'include'))
41	include_dirs.append(os.path.join(torch_npu_dir, 'include', 'third_party', 'acl', 'inc'))	41	include_dirs.append(os.path.join(torch_npu_dir, 'include', 'third_party', 'acl', 'inc'))
42	include_dirs.append(os.path.join(torch_npu_dir, 'include', 'third_party', 'hccl', 'inc'))	42	include_dirs.append(os.path.join(torch_npu_dir, 'include', 'third_party', 'hccl', 'inc'))
43	include_dirs += TorchExtension.include_paths()	43	include_dirs += TorchExtension.include_paths()
44	kwargs['include_dirs'] = include_dirs	44	kwargs['include_dirs'] = include_dirs
45		45
46	library_dirs = kwargs.get('library_dirs', [])	46	library_dirs = kwargs.get('library_dirs', [])
47	library_dirs.append(os.path.join(torch_npu_dir, 'lib'))	47	library_dirs.append(os.path.join(torch_npu_dir, 'lib'))
48	library_dirs += TorchExtension.library_paths()	48	library_dirs += TorchExtension.library_paths()
49	kwargs['library_dirs'] = library_dirs	49	kwargs['library_dirs'] = library_dirs
50		50
51	libraries = kwargs.get('libraries', [])	51	libraries = kwargs.get('libraries', [])
52	libraries.append('c10')	52	libraries.append('c10')
53	libraries.append('torch')	53	libraries.append('torch')
54	libraries.append('torch_cpu')	54	libraries.append('torch_cpu')
55	libraries.append('torch_python')	55	libraries.append('torch_python')
56	libraries.append('torch_npu')	56	libraries.append('torch_npu')
57	kwargs['libraries'] = libraries	57	kwargs['libraries'] = libraries
58		58
59	kwargs['language'] = 'c++'	59	kwargs['language'] = 'c++'
60	return setuptools.Extension(name, sources, args, *kwargs)	60	return setuptools.Extension(name, sources, args, *kwargs)

Mtorch_npu/utils/flops_count.py+0-1

	@@ -38,4 +38,3 @@ class FlopsCounter(_FlopsCounter):
38	super().__init__()	38	super().__init__()
39	warnings.warn("torch_npu.utils.flops_count.FlopsCounter() will be deprecated. "	39	warnings.warn("torch_npu.utils.flops_count.FlopsCounter() will be deprecated. "
40	"If necessary, please use torch_npu.utils.FlopsCounter().", FutureWarning)	40	"If necessary, please use torch_npu.utils.FlopsCounter().", FutureWarning)
41

Mtorch_npu/utils/syncbatchnorm.py+103-103

	@@ -1,103 +1,103 @@
1	import torch	1	import torch
2	import torch.distributed as dist	2	import torch.distributed as dist
3	from torch.autograd.function import Function	3	from torch.autograd.function import Function
4		4
5	import torch_npu	5	import torch_npu
6	from torch_npu.utils._error_code import ErrCode, ops_error	6	from torch_npu.utils._error_code import ErrCode, ops_error
7		7
8		8
9	__all__ = ["SyncBatchNorm"]	9	__all__ = ["SyncBatchNorm"]
10		10
11		11
12	class SyncBatchNorm(Function):	12	class SyncBatchNorm(Function):
13		13
14	@staticmethod	14	@staticmethod
15	def forward(self, input_tensor, weight, bias, running_mean, running_var, eps, momentum, process_group, world_size):	15	def forward(self, input_tensor, weight, bias, running_mean, running_var, eps, momentum, process_group, world_size):
16	input_tensor = input_tensor.contiguous()	16	input_tensor = input_tensor.contiguous()
17	input_shape = input_tensor.shape	17	input_shape = input_tensor.shape
18	input_tensor_ = input_tensor.reshape(input_shape[0], input_shape[1], 1, -1)	18	input_tensor_ = input_tensor.reshape(input_shape[0], input_shape[1], 1, -1)
19	# calculate sum/sum_square for input.	19	# calculate sum/sum_square for input.
20	sum_val, sum_square_val = torch_npu.batch_norm_reduce(input_tensor_, eps)	20	sum_val, sum_square_val = torch_npu.batch_norm_reduce(input_tensor_, eps)
21		21
22	count = torch.full((1,),	22	count = torch.full((1,),
23	input_tensor.numel() // input_tensor.size(1),	23	input_tensor.numel() // input_tensor.size(1),
24	dtype=sum_val.dtype,	24	dtype=sum_val.dtype,
25	device=sum_val.device)	25	device=sum_val.device)
26		26
27	num_channels = input_tensor.shape[1]	27	num_channels = input_tensor.shape[1]
28	# C, C, 1 -> (2C + 1)	28	# C, C, 1 -> (2C + 1)
29	combined = torch.cat([sum_val, sum_square_val, count], dim=0)	29	combined = torch.cat([sum_val, sum_square_val, count], dim=0)
30	# world_size * (2C + 1)	30	# world_size * (2C + 1)
31	combined_list = [torch.empty_like(combined) for k in range(world_size)]	31	combined_list = [torch.empty_like(combined) for k in range(world_size)]
32	dist.all_gather(combined_list, combined, process_group, async_op=False)	32	dist.all_gather(combined_list, combined, process_group, async_op=False)
33	combined = torch.stack(combined_list, dim=0)	33	combined = torch.stack(combined_list, dim=0)
34	# world_size * (2C + 1) -> world_size * C, world_size * C, world_size * 1	34	# world_size * (2C + 1) -> world_size * C, world_size * C, world_size * 1
35	sum_all, square_sum_all, count_all = torch.split(combined, num_channels, dim=1)	35	sum_all, square_sum_all, count_all = torch.split(combined, num_channels, dim=1)
36		36
37	size = count_all.view(-1).sum()	37	size = count_all.view(-1).sum()
38	if size == 1:	38	if size == 1:
39	raise ValueError('Expected more than 1 value per channel when training, got input size {}'.format(size) +	39	raise ValueError('Expected more than 1 value per channel when training, got input size {}'.format(size) +
40	ops_error(ErrCode.VALUE))	40	ops_error(ErrCode.VALUE))
41		41
42	# calculate global mean & invstd	42	# calculate global mean & invstd
43	mean, invstd = torch_npu.batch_norm_gather_stats_update(input_tensor,	43	mean, invstd = torch_npu.batch_norm_gather_stats_update(input_tensor,
44	sum_all,	44	sum_all,
45	square_sum_all,	45	square_sum_all,
46	running_mean,	46	running_mean,
47	running_var,	47	running_var,
48	momentum,	48	momentum,
49	eps,	49	eps,
50	count_all.view(-1))	50	count_all.view(-1))
51		51
52	self.save_for_backward(input_tensor, weight, mean, invstd, count_all)	52	self.save_for_backward(input_tensor, weight, mean, invstd, count_all)
53	self.process_group = process_group	53	self.process_group = process_group
54		54
55	# apply element-wise normalization	55	# apply element-wise normalization
56	out = torch.batch_norm_elemt(input_tensor, weight, bias, mean, invstd, eps)	56	out = torch.batch_norm_elemt(input_tensor, weight, bias, mean, invstd, eps)
57	return out	57	return out
58		58
59	@staticmethod	59	@staticmethod
60	def backward(self, grad_output):	60	def backward(self, grad_output):
61	if not grad_output.is_contiguous(memory_format=torch.channels_last):	61	if not grad_output.is_contiguous(memory_format=torch.channels_last):
62	grad_output = grad_output.contiguous()	62	grad_output = grad_output.contiguous()
63	saved_input, weight, mean, invstd, count_tensor = self.saved_tensors	63	saved_input, weight, mean, invstd, count_tensor = self.saved_tensors
64	grad_input = grad_weight = grad_bias = None	64	grad_input = grad_weight = grad_bias = None
65	process_group = self.process_group	65	process_group = self.process_group
66		66
67	# calculate local stats as well as grad_weight / grad_bias	67	# calculate local stats as well as grad_weight / grad_bias
68	sum_dy, sum_dy_xmu, grad_weight, grad_bias = torch.batch_norm_backward_reduce(grad_output,	68	sum_dy, sum_dy_xmu, grad_weight, grad_bias = torch.batch_norm_backward_reduce(grad_output,
69	saved_input,	69	saved_input,
70	mean,	70	mean,
71	invstd,	71	invstd,
72	weight,	72	weight,
73	self.needs_input_grad[0],	73	self.needs_input_grad[0],
74	self.needs_input_grad[1],	74	self.needs_input_grad[1],
75	self.needs_input_grad[2])	75	self.needs_input_grad[2])
76		76
77	if self.needs_input_grad[0]:	77	if self.needs_input_grad[0]:
78	# synchronizing stats used to calculate input gradient.	78	# synchronizing stats used to calculate input gradient.
79	num_channels = sum_dy.shape[0]	79	num_channels = sum_dy.shape[0]
80	combined = torch.cat([sum_dy, sum_dy_xmu], dim=0)	80	combined = torch.cat([sum_dy, sum_dy_xmu], dim=0)
81	torch.distributed.all_reduce(	81	torch.distributed.all_reduce(
82	combined, torch.distributed.ReduceOp.SUM, process_group, async_op=False)	82	combined, torch.distributed.ReduceOp.SUM, process_group, async_op=False)
83	sum_dy, sum_dy_xmu = torch.split(combined, num_channels)	83	sum_dy, sum_dy_xmu = torch.split(combined, num_channels)
84		84
85	# backward pass for gradient calculation	85	# backward pass for gradient calculation
86	grad_input = torch.batch_norm_backward_elemt(grad_output,	86	grad_input = torch.batch_norm_backward_elemt(grad_output,
87	saved_input,	87	saved_input,
88	mean,	88	mean,
89	invstd,	89	invstd,
90	weight,	90	weight,
91	sum_dy,	91	sum_dy,
92	sum_dy_xmu,	92	sum_dy_xmu,
93	count_tensor)	93	count_tensor)
94		94
95	# synchronizing of grad_weight / grad_bias is not needed as distributed	95	# synchronizing of grad_weight / grad_bias is not needed as distributed
96	# training would handle all reduce.	96	# training would handle all reduce.
97	if weight is None or not self.needs_input_grad[1]:	97	if weight is None or not self.needs_input_grad[1]:
98	grad_weight = None	98	grad_weight = None
99		99
100	if weight is None or not self.needs_input_grad[2]:	100	if weight is None or not self.needs_input_grad[2]:
101	grad_bias = None	101	grad_bias = None
102		102
103	return grad_input, grad_weight, grad_bias, None, None, None, None, None, None	103	return grad_input, grad_weight, grad_bias, None, None, None, None, None, None

Mtorch_npu/utils/tensor_methods.py+88-88

	@@ -1,88 +1,88 @@
1	from functools import wraps	1	from functools import wraps
2		2
3	import torch	3	import torch
4		4
5	import torch_npu	5	import torch_npu
6	from torch_npu.utils._error_code import ErrCode, pta_error	6	from torch_npu.utils._error_code import ErrCode, pta_error
7	from torch_npu.utils.storage import _reduce_ex	7	from torch_npu.utils.storage import _reduce_ex
8		8
9		9
10	__all__ = []	10	__all__ = []
11		11
12		12
13	def _npu(self, args, *kwargs):	13	def _npu(self, args, *kwargs):
14	return torch_npu._C.npu(self, args, *kwargs)	14	return torch_npu._C.npu(self, args, *kwargs)
15		15
16		16
17	@property	17	@property
18	def _is_npu(self):	18	def _is_npu(self):
19	return torch_npu._C.is_npu(self)	19	return torch_npu._C.is_npu(self)
20		20
21		21
22	class _NPUTensortypeCache(object):	22	class _NPUTensortypeCache(object):
23	init = False	23	init = False
24	tensortype_list = []	24	tensortype_list = []
25	tensortype_dict = {}	25	tensortype_dict = {}
26		26
27	@classmethod	27	@classmethod
28	def tensortype_list_dict_init(cls):	28	def tensortype_list_dict_init(cls):
29	if not cls.init:	29	if not cls.init:
30	cls.tensortype_list += [	30	cls.tensortype_list += [
31	torch_npu.npu.BoolTensor,	31	torch_npu.npu.BoolTensor,
32	torch_npu.npu.ByteTensor,	32	torch_npu.npu.ByteTensor,
33	torch_npu.npu.CharTensor,	33	torch_npu.npu.CharTensor,
34	torch_npu.npu.DoubleTensor,	34	torch_npu.npu.DoubleTensor,
35	torch_npu.npu.FloatTensor,	35	torch_npu.npu.FloatTensor,
36	torch_npu.npu.HalfTensor,	36	torch_npu.npu.HalfTensor,
37	torch_npu.npu.IntTensor,	37	torch_npu.npu.IntTensor,
38	torch_npu.npu.LongTensor,	38	torch_npu.npu.LongTensor,
39	torch_npu.npu.ShortTensor,	39	torch_npu.npu.ShortTensor,
40	torch_npu.npu.BFloat16Tensor,	40	torch_npu.npu.BFloat16Tensor,
41	]	41	]
42		42
43	cls.tensortype_str_list = [	43	cls.tensortype_str_list = [
44	"torch_npu.npu.BoolTensor",	44	"torch_npu.npu.BoolTensor",
45	"torch_npu.npu.ByteTensor",	45	"torch_npu.npu.ByteTensor",
46	"torch_npu.npu.CharTensor",	46	"torch_npu.npu.CharTensor",
47	"torch_npu.npu.DoubleTensor",	47	"torch_npu.npu.DoubleTensor",
48	"torch_npu.npu.FloatTensor",	48	"torch_npu.npu.FloatTensor",
49	"torch_npu.npu.HalfTensor",	49	"torch_npu.npu.HalfTensor",
50	"torch_npu.npu.IntTensor",	50	"torch_npu.npu.IntTensor",
51	"torch_npu.npu.LongTensor",	51	"torch_npu.npu.LongTensor",
52	"torch_npu.npu.ShortTensor",	52	"torch_npu.npu.ShortTensor",
53	"torch_npu.npu.BFloat16Tensor",	53	"torch_npu.npu.BFloat16Tensor",
54	]	54	]
55		55
56	for tensortype, tensortype_str in zip(cls.tensortype_list, cls.tensortype_str_list):	56	for tensortype, tensortype_str in zip(cls.tensortype_list, cls.tensortype_str_list):
57	cls.tensortype_dict[tensortype_str] = tensortype	57	cls.tensortype_dict[tensortype_str] = tensortype
58	cls.tensortype_dict[tensortype_str.replace('torch_npu.', 'torch.')] = tensortype	58	cls.tensortype_dict[tensortype_str.replace('torch_npu.', 'torch.')] = tensortype
59		59
60	cls.init = True	60	cls.init = True
61		61
62	@classmethod	62	@classmethod
63	def get_tensortype_list(cls):	63	def get_tensortype_list(cls):
64	return cls.tensortype_list	64	return cls.tensortype_list
65		65
66	@classmethod	66	@classmethod
67	def get_tensortype_dict(cls):	67	def get_tensortype_dict(cls):
68	return cls.tensortype_dict	68	return cls.tensortype_dict
69		69
70		70
71	def _npu_type(self, dtype=None, non_blocking=False, **kwargs):	71	def _npu_type(self, dtype=None, non_blocking=False, **kwargs):
72	if dtype is None:	72	if dtype is None:
73	return self.type_raw(dtype, non_blocking, **kwargs)	73	return self.type_raw(dtype, non_blocking, **kwargs)
74		74
75	_NPUTensortypeCache.tensortype_list_dict_init()	75	_NPUTensortypeCache.tensortype_list_dict_init()
76	if isinstance(dtype, str) and dtype in _NPUTensortypeCache.get_tensortype_dict():	76	if isinstance(dtype, str) and dtype in _NPUTensortypeCache.get_tensortype_dict():
77	tensortype_class = _NPUTensortypeCache.get_tensortype_dict()[dtype]	77	tensortype_class = _NPUTensortypeCache.get_tensortype_dict()[dtype]
78	return self.to(dtype=tensortype_class.dtype, device='npu', non_blocking=non_blocking)	78	return self.to(dtype=tensortype_class.dtype, device='npu', non_blocking=non_blocking)
79	elif dtype in _NPUTensortypeCache.get_tensortype_list():	79	elif dtype in _NPUTensortypeCache.get_tensortype_list():
80	return self.to(dtype=dtype.dtype, device='npu', non_blocking=non_blocking)	80	return self.to(dtype=dtype.dtype, device='npu', non_blocking=non_blocking)
81	else:	81	else:
82	return self.type_raw(dtype, non_blocking, **kwargs)	82	return self.type_raw(dtype, non_blocking, **kwargs)
83		83
84		84
85	def _add_tensor_methods():	85	def _add_tensor_methods():
86	torch.Tensor.type_raw = torch.Tensor.type	86	torch.Tensor.type_raw = torch.Tensor.type
87	torch.Tensor.type = _npu_type	87	torch.Tensor.type = _npu_type
88	torch.Tensor.__reduce_ex__ = _reduce_ex	88	torch.Tensor.__reduce_ex__ = _reduce_ex

	@@ -1,190 +1,190 @@
1	OPEN SOURCE SOFTWARE NOTICE	1	OPEN SOURCE SOFTWARE NOTICE
2	Please note we provide an open source software notice along with this product and/or this product firmware (in the following just "this product"). The open source software licenses are granted by the respective right holders. And the open source licenses prevail all other license information with regard to the respective open source software contained in the product, including but not limited to End User Software Licensing Agreement. This notice is provided on behalf of Huawei Technologies Co. Ltd. and any of its local subsidiaries which may have provided this product to you in your local country.	2	Please note we provide an open source software notice along with this product and/or this product firmware (in the following just "this product"). The open source software licenses are granted by the respective right holders. And the open source licenses prevail all other license information with regard to the respective open source software contained in the product, including but not limited to End User Software Licensing Agreement. This notice is provided on behalf of Huawei Technologies Co. Ltd. and any of its local subsidiaries which may have provided this product to you in your local country.
3		3
4	Warranty Disclaimer	4	Warranty Disclaimer
5	THE OPEN SOURCE SOFTWARE IN THIS PRODUCT IS DISTRIBUTED IN THE HOPE THAT IT WILL BE USEFUL, BUT WITHOUT ANY WARRANTY, WITHOUT EVEN THE IMPLIED WARRANTY OF MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE. SEE THE APPLICABLE LICENSES FOR MORE DETAILS.	5	THE OPEN SOURCE SOFTWARE IN THIS PRODUCT IS DISTRIBUTED IN THE HOPE THAT IT WILL BE USEFUL, BUT WITHOUT ANY WARRANTY, WITHOUT EVEN THE IMPLIED WARRANTY OF MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE. SEE THE APPLICABLE LICENSES FOR MORE DETAILS.
6		6
7	Copyright Notice and License Texts	7	Copyright Notice and License Texts
8	Software: pytorch v2.6.0	8	Software: pytorch v2.6.0
9	Copyright notice:	9	Copyright notice:
10	Copyright (c) Advanced Micro Devices, Inc.	10	Copyright (c) Advanced Micro Devices, Inc.
11		11
12	Copyright (c) Microsoft Corporation	12	Copyright (c) Microsoft Corporation
13		13
14	Copyright (c) Bjorn Fahller	14	Copyright (c) Bjorn Fahller
15		15
16	Copyright (c) 2001-2014 Python Software Foundation All Rights Reserved	16	Copyright (c) 2001-2014 Python Software Foundation All Rights Reserved
17		17
18	Copyright (c) 2011-2013 NYU	18	Copyright (c) 2011-2013 NYU
19		19
20	Copyright (c) 1995-2011 by Fredrik Lundh	20	Copyright (c) 1995-2011 by Fredrik Lundh
21		21
22	Copyright (c) Edward Z. Yang ezyang@mit.edu	22	Copyright (c) Edward Z. Yang ezyang@mit.edu
23		23
24	Copyright (c) 2014- Facebook, Inc	24	Copyright (c) 2014- Facebook, Inc
25		25
26	Copyright (c) 2017 The Android Open Source Project	26	Copyright (c) 2017 The Android Open Source Project
27		27
28	Copyright Python Software Foundation	28	Copyright Python Software Foundation
29		29
30	Copyright (c) 2012 Massachusetts Institute of Technology	30	Copyright (c) 2012 Massachusetts Institute of Technology
31		31
32	Copyright (c) 2018 Alex Rogozhnikov	32	Copyright (c) 2018 Alex Rogozhnikov
33		33
34	Copyright (c) 2007-2009 Scientific Computing and Imaging Institute, University of Utah	34	Copyright (c) 2007-2009 Scientific Computing and Imaging Institute, University of Utah
35		35
36	Copyright (c) 2006 Idiap Research Institute	36	Copyright (c) 2006 Idiap Research Institute
37		37
38	Copyright (c) 2011-2021, NVIDIA CORPORATION. All rights reserved	38	Copyright (c) 2011-2021, NVIDIA CORPORATION. All rights reserved
39		39
40	Copyright (c) 2015 Yangqing Jia All rights reserved	40	Copyright (c) 2015 Yangqing Jia All rights reserved
41		41
42	Copyright (c) Meta Platforms, Inc.	42	Copyright (c) Meta Platforms, Inc.
43		43
44	Copyright 2023-present Facebook. All Rights Reserved	44	Copyright 2023-present Facebook. All Rights Reserved
45		45
46	Copyright (c) 2022 Apple Inc.	46	Copyright (c) 2022 Apple Inc.
47		47
48	Copyright (c) 2005-2017, NumPy Developers. All rights reserved	48	Copyright (c) 2005-2017, NumPy Developers. All rights reserved
49		49
50	Copyright (c) 1997 - 2002, Makoto Matsumoto and Takuji Nishimura, All rights reserved	50	Copyright (c) 1997 - 2002, Makoto Matsumoto and Takuji Nishimura, All rights reserved
51		51
52	Copyright (c) 2014, The Regents	52	Copyright (c) 2014, The Regents
53		53
54	Copyright (c) 2005-2010 ActiveState Software Inc.	54	Copyright (c) 2005-2010 ActiveState Software Inc.
55		55
56	Copyright Malte Skarupke 2017	56	Copyright Malte Skarupke 2017
57		57
58	Copyright (c) 2006-2010 NEC Laboratories America (Ronan Collobert, Leon Bottou, Iain Melvin, Jason Weston)	58	Copyright (c) 2006-2010 NEC Laboratories America (Ronan Collobert, Leon Bottou, Iain Melvin, Jason Weston)
59		59
60	Copyright 2005, Google Inc. All rights reserved	60	Copyright 2005, Google Inc. All rights reserved
61		61
62	Copyright (c) Meta Platforms, Inc. and affiliates	62	Copyright (c) Meta Platforms, Inc. and affiliates
63		63
64	Copyright (c) 2023, Advanced Micro Devices, Inc.	64	Copyright (c) 2023, Advanced Micro Devices, Inc.
65		65
66	Copyright (c) 2022, Tri Dao	66	Copyright (c) 2022, Tri Dao
67		67
68	Copyright (c) 2005-2023 NVIDIA Corporation Built	68	Copyright (c) 2005-2023 NVIDIA Corporation Built
69		69
70	Copyright (c) 2001-2002 Enthought, Inc. 2003-2019, SciPy Developers. All rights reserved	70	Copyright (c) 2001-2002 Enthought, Inc. 2003-2019, SciPy Developers. All rights reserved
71		71
72	Copyright 2008 Google Inc. All rights reserved	72	Copyright 2008 Google Inc. All rights reserved
73		73
74	Copyright (c) 2021, 2023-2024 Arm Limited	74	Copyright (c) 2021, 2023-2024 Arm Limited
75		75
76	Copyright (c) 2003-2017 Josef Weidendorfer. All rights reserved	76	Copyright (c) 2003-2017 Josef Weidendorfer. All rights reserved
77		77
78	Copyright (c) 1997-2011 by Secret Labs AB	78	Copyright (c) 1997-2011 by Secret Labs AB
79		79
80	Copyright (c) 2016- Facebook, Inc	80	Copyright (c) 2016- Facebook, Inc
81		81
82	Copyright (c) 2014 Matthew Rocklin	82	Copyright (c) 2014 Matthew Rocklin
83		83
84	Copyright (c) 2005-2022 NVIDIA Corporation Built	84	Copyright (c) 2005-2022 NVIDIA Corporation Built
85		85
86	Copyright (c) Facebook, Inc.	86	Copyright (c) Facebook, Inc.
87		87
88	Copyright 2019-2020 Kakao Brain	88	Copyright 2019-2020 Kakao Brain
89		89
90	Copyright (c) 2000-2017 Julian Seward. All rights reserved	90	Copyright (c) 2000-2017 Julian Seward. All rights reserved
91		91
92	Copyright (c) 2005-2020 Rich Felker	92	Copyright (c) 2005-2020 Rich Felker
93		93
94	Copyright (c) 2008 - 2009 NVIDIA Corporation. All rights reserved	94	Copyright (c) 2008 - 2009 NVIDIA Corporation. All rights reserved
95		95
96	Copyright 2010-2014 Rich Geldreich and Tenacious Software LLC	96	Copyright 2010-2014 Rich Geldreich and Tenacious Software LLC
97		97
98	Copyright (c) 2016 Facebook Inc.	98	Copyright (c) 2016 Facebook Inc.
99		99
100	Copyright (c) 2001-2004 Idiap Research Institute (Ronan Collobert, Samy Bengio, Johnny Mariethoz)	100	Copyright (c) 2001-2004 Idiap Research Institute (Ronan Collobert, Samy Bengio, Johnny Mariethoz)
101		101
102	Copyright 2010-2014 Rich Geldreich and Tenacious Software LLC All Rights Reserved	102	Copyright 2010-2014 Rich Geldreich and Tenacious Software LLC All Rights Reserved
103		103
104	Copyright (c) 2012-2014 Deepmind Technologies	104	Copyright (c) 2012-2014 Deepmind Technologies
105		105
106	Copyright (c) 2012 Giovanni Garberoglio Interdisciplinary Laboratory	106	Copyright (c) 2012 Giovanni Garberoglio Interdisciplinary Laboratory
107		107
108	Copyright (c) 2024, Tri Dao	108	Copyright (c) 2024, Tri Dao
109		109
110	Copyright (c) Donald Stufft and individual contributors. All rights reserved	110	Copyright (c) Donald Stufft and individual contributors. All rights reserved
111		111
112	Copyright (c) 2018, Steven Moshier All rights reserved	112	Copyright (c) 2018, Steven Moshier All rights reserved
113		113
114	Copyright (c) 2015 Google Inc. All rights reserved	114	Copyright (c) 2015 Google Inc. All rights reserved
115		115
116	Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved	116	Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved
117		117
118	Copyright (c) 2010-2022 by Alex Clark and contributors	118	Copyright (c) 2010-2022 by Alex Clark and contributors
119		119
120	Copyright 2015 Google Inc. All Rights Reserved	120	Copyright 2015 Google Inc. All Rights Reserved
121		121
122	Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved	122	Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved
123		123
124	Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved	124	Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved
125		125
126	Copyright (c) 2016 manylinux	126	Copyright (c) 2016 manylinux
127		127
128	Copyright (c) 2017 - 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved	128	Copyright (c) 2017 - 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved
129		129
130	Copyright (c) 2011-2012 NEC Laboratories America (Koray Kavukcuoglu)	130	Copyright (c) 2011-2012 NEC Laboratories America (Koray Kavukcuoglu)
131		131
132	Copyright 2013-2014 RAD Game	132	Copyright 2013-2014 RAD Game
133		133
134	Copyright (c) 2011-2019 Stephan Brumme. All rights reserved	134	Copyright (c) 2011-2019 Stephan Brumme. All rights reserved
135		135
136	Copyright (c) 2018 MathInf GmbH, Thomas Viehmann	136	Copyright (c) 2018 MathInf GmbH, Thomas Viehmann
137		137
138	Copyright (c) 2013 Eddy Petrisor	138	Copyright (c) 2013 Eddy Petrisor
139		139
140	Copyright (c) 2023-2024 The ggml	140	Copyright (c) 2023-2024 The ggml
141		141
142	Copyright (c) 2001, 2002, 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010, 2011, 2012, 2013, 2014, 2015, 2016, 2017, 2018, 2019, 2020 Python Software Foundation All Rights Reserved	142	Copyright (c) 2001, 2002, 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010, 2011, 2012, 2013, 2014, 2015, 2016, 2017, 2018, 2019, 2020 Python Software Foundation All Rights Reserved
143		143
144	Copyright 2004-present Facebook. All Rights Reserved	144	Copyright 2004-present Facebook. All Rights Reserved
145		145
146	Copyright (c) 2010 ActiveState Software Inc.	146	Copyright (c) 2010 ActiveState Software Inc.
147		147
148	Copyright (c) 2006 The Android Open Source Project	148	Copyright (c) 2006 The Android Open Source Project
149		149
150	(c) Meta Platforms, Inc. and affiliates. All rights reserved	150	(c) Meta Platforms, Inc. and affiliates. All rights reserved
151		151
152	Copyright (c) 2023 Apple Inc.	152	Copyright (c) 2023 Apple Inc.
153		153
154	Copyright (c) Microsoft Corporation. All rights reserved	154	Copyright (c) Microsoft Corporation. All rights reserved
155		155
156	Copyright 2015 The TensorFlow Authors. All Rights Reserved	156	Copyright 2015 The TensorFlow Authors. All Rights Reserved
157		157
158	Copyright (c) 2023, Tri Dao	158	Copyright (c) 2023, Tri Dao
159		159
160	Copyright 2022 Cruise LLC	160	Copyright 2022 Cruise LLC
161		161
162	Copyright (c) Meta Platforms, Inc. and affiliates. All rights reserved	162	Copyright (c) Meta Platforms, Inc. and affiliates. All rights reserved
163		163
164	Copyright (c) 2022 Cruise LLC. All rights reserved	164	Copyright (c) 2022 Cruise LLC. All rights reserved
165		165
166	Copyright (c) 2016-present, Facebook, Inc.	166	Copyright (c) 2016-present, Facebook, Inc.
167		167
168	(c) Copyright John Maddock 2006	168	(c) Copyright John Maddock 2006
169		169
170	Copyright (c) 2011-2014 Idiap Research Institute	170	Copyright (c) 2011-2014 Idiap Research Institute
171		171
172	Copyright (c) 2014 Indiana University All rights reserved	172	Copyright (c) 2014 Indiana University All rights reserved
173		173
174	copyright 2019 The TensorFlow Authors	174	copyright 2019 The TensorFlow Authors
175		175
176	Copyright (c) 2016-present, Facebook Inc. All rights reserved	176	Copyright (c) 2016-present, Facebook Inc. All rights reserved
177		177
178	License: BSD 3-Clause License	178	License: BSD 3-Clause License
179	Copyright (c) ,	179	Copyright (c) ,
180	All rights reserved.	180	All rights reserved.
181		181
182	1. Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:	182	1. Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:
183		183
184	2. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.	184	2. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.
185		185
186	3. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.	186	3. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.
187		187
188	Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission.	188	Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission.
189		189
190	THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.	190	THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

	@@ -1,92 +1,92 @@
1	import torch	1	import torch
2	import torch.fx as fx	2	import torch.fx as fx
3	from torch.fx.passes.shape_prop import ShapeProp	3	from torch.fx.passes.shape_prop import ShapeProp
4	from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests	4	from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
5	from testutils import TestUtils	5	from testutils import TestUtils
6	import torch_npu	6	import torch_npu
7	import torch_npu._inductor	7	import torch_npu._inductor
8	from torch_npu._inductor.fx_passes.utils.fx_pass_level import FxPassLevel	8	from torch_npu._inductor.fx_passes.utils.fx_pass_level import FxPassLevel
9	from torch_npu._inductor.fx_passes.utils.fx_pass_level import PassType	9	from torch_npu._inductor.fx_passes.utils.fx_pass_level import PassType
10		10
11		11
12	class TestFxPassLevel(TestUtils):	12	class TestFxPassLevel(TestUtils):
13		13
14	def test_comparison_between_members(self):	14	def test_comparison_between_members(self):
15	"""枚举成员之间的大小比较"""	15	"""枚举成员之间的大小比较"""
16	assert FxPassLevel.LEVEL1 < FxPassLevel.LEVEL2	16	assert FxPassLevel.LEVEL1 < FxPassLevel.LEVEL2
17	assert FxPassLevel.LEVEL1 <= FxPassLevel.LEVEL2	17	assert FxPassLevel.LEVEL1 <= FxPassLevel.LEVEL2
18	assert FxPassLevel.LEVEL2 > FxPassLevel.LEVEL1	18	assert FxPassLevel.LEVEL2 > FxPassLevel.LEVEL1
19	assert FxPassLevel.LEVEL2 >= FxPassLevel.LEVEL1	19	assert FxPassLevel.LEVEL2 >= FxPassLevel.LEVEL1
20	assert FxPassLevel.LEVEL2 == FxPassLevel.LEVEL2	20	assert FxPassLevel.LEVEL2 == FxPassLevel.LEVEL2
21	assert FxPassLevel.LEVEL1 != FxPassLevel.LEVEL3	21	assert FxPassLevel.LEVEL1 != FxPassLevel.LEVEL3
22		22
23	# 严格顺序	23	# 严格顺序
24	assert FxPassLevel.LEVEL1 < FxPassLevel.LEVEL2 < FxPassLevel.LEVEL3	24	assert FxPassLevel.LEVEL1 < FxPassLevel.LEVEL2 < FxPassLevel.LEVEL3
25		25
26		26
27	def test_sorting_and_min_max(self):	27	def test_sorting_and_min_max(self):
28	"""支持排序、min、max 等操作"""	28	"""支持排序、min、max 等操作"""
29	levels = [FxPassLevel.LEVEL3, FxPassLevel.LEVEL1, FxPassLevel.LEVEL2]	29	levels = [FxPassLevel.LEVEL3, FxPassLevel.LEVEL1, FxPassLevel.LEVEL2]
30	sorted_levels = sorted(levels)	30	sorted_levels = sorted(levels)
31	assert sorted_levels == [FxPassLevel.LEVEL1, FxPassLevel.LEVEL2, FxPassLevel.LEVEL3]	31	assert sorted_levels == [FxPassLevel.LEVEL1, FxPassLevel.LEVEL2, FxPassLevel.LEVEL3]
32		32
33	assert min(levels) == FxPassLevel.LEVEL1	33	assert min(levels) == FxPassLevel.LEVEL1
34	assert max(levels) == FxPassLevel.LEVEL3	34	assert max(levels) == FxPassLevel.LEVEL3
35		35
36		36
37	def test_hash_consistency(self):	37	def test_hash_consistency(self):
38	"""哈希值一致性"""	38	"""哈希值一致性"""
39	assert hash(FxPassLevel.LEVEL1) == hash(1)	39	assert hash(FxPassLevel.LEVEL1) == hash(1)
40	assert hash(FxPassLevel.LEVEL2) == hash(2)	40	assert hash(FxPassLevel.LEVEL2) == hash(2)
41	assert hash(FxPassLevel.LEVEL3) == hash(3)	41	assert hash(FxPassLevel.LEVEL3) == hash(3)
42		42
43	# 相同值哈希相同	43	# 相同值哈希相同
44	assert hash(FxPassLevel.LEVEL1) == hash(FxPassLevel.LEVEL1)	44	assert hash(FxPassLevel.LEVEL1) == hash(FxPassLevel.LEVEL1)
45		45
46	# 不同值哈希不同	46	# 不同值哈希不同
47	assert hash(FxPassLevel.LEVEL1) != hash(FxPassLevel.LEVEL2)	47	assert hash(FxPassLevel.LEVEL1) != hash(FxPassLevel.LEVEL2)
48		48
49	# 可用于集合	49	# 可用于集合
50	s = {FxPassLevel.LEVEL1, FxPassLevel.LEVEL2}	50	s = {FxPassLevel.LEVEL1, FxPassLevel.LEVEL2}
51	assert FxPassLevel.LEVEL1 in s	51	assert FxPassLevel.LEVEL1 in s
52	assert FxPassLevel.LEVEL3 not in s	52	assert FxPassLevel.LEVEL3 not in s
53		53
54		54
55	class TestPassType(TestUtils):	55	class TestPassType(TestUtils):
56		56
57	def test_comparison_between_members(self):	57	def test_comparison_between_members(self):
58	"""枚举成员之间的大小比较"""	58	"""枚举成员之间的大小比较"""
59	assert PassType.PRE < PassType.POST	59	assert PassType.PRE < PassType.POST
60	assert PassType.PRE <= PassType.POST	60	assert PassType.PRE <= PassType.POST
61	assert PassType.POST > PassType.PRE	61	assert PassType.POST > PassType.PRE
62	assert PassType.POST >= PassType.PRE	62	assert PassType.POST >= PassType.PRE
63	assert PassType.PRE == PassType.PRE	63	assert PassType.PRE == PassType.PRE
64	assert PassType.PRE != PassType.POST	64	assert PassType.PRE != PassType.POST
65		65
66		66
67	def test_sorting_and_min_max(self):	67	def test_sorting_and_min_max(self):
68	"""支持排序、min、max"""	68	"""支持排序、min、max"""
69	types = [PassType.POST, PassType.PRE]	69	types = [PassType.POST, PassType.PRE]
70	sorted_types = sorted(types)	70	sorted_types = sorted(types)
71	assert sorted_types == [PassType.PRE, PassType.POST]	71	assert sorted_types == [PassType.PRE, PassType.POST]
72		72
73	assert min(types) == PassType.PRE	73	assert min(types) == PassType.PRE
74	assert max(types) == PassType.POST	74	assert max(types) == PassType.POST
75		75
76		76
77	def test_hash_consistency(self):	77	def test_hash_consistency(self):
78	"""哈希值一致性"""	78	"""哈希值一致性"""
79	assert hash(PassType.PRE) == hash(1)	79	assert hash(PassType.PRE) == hash(1)
80	assert hash(PassType.POST) == hash(2)	80	assert hash(PassType.POST) == hash(2)
81		81
82	s = {PassType.PRE, PassType.POST}	82	s = {PassType.PRE, PassType.POST}
83	assert PassType.PRE in s	83	assert PassType.PRE in s
84	assert PassType.PRE in s # 重复添加不影响	84	assert PassType.PRE in s # 重复添加不影响
85		85
86		86
87	instantiate_parametrized_tests(TestFxPassLevel)	87	instantiate_parametrized_tests(TestFxPassLevel)
88	instantiate_parametrized_tests(TestPassType)	88	instantiate_parametrized_tests(TestPassType)
89		89
90		90
91	if __name__ == "__main__":	91	if __name__ == "__main__":
92	run_tests()	92	run_tests()

	@@ -1,97 +1,97 @@
1	import torch	1	import torch
2	from torch.testing._internal.common_utils import run_tests	2	from torch.testing._internal.common_utils import run_tests
3	from testutils import TestUtils	3	from testutils import TestUtils
4		4
5	from torch_npu._inductor.codegen.tile_generator import TileGenerator, aligned_numel_32byte	5	from torch_npu._inductor.codegen.tile_generator import TileGenerator, aligned_numel_32byte
6	from torch_npu._inductor.codegen.triton_utils import NPUKernelType	6	from torch_npu._inductor.codegen.triton_utils import NPUKernelType
7		7
8		8
9	class TestTileGenerator32ByteAlignment(TestUtils):	9	class TestTileGenerator32ByteAlignment(TestUtils):
10	"""	10	"""
11	Test class for TileGenerator 32Byte alignment functionality	11	Test class for TileGenerator 32Byte alignment functionality
12	"""	12	"""
13		13
14	def test_32byte_alignment(self):	14	def test_32byte_alignment(self):
15	"""	15	"""
16	Test that block_size and sub_block_size are 32Byte aligned	16	Test that block_size and sub_block_size are 32Byte aligned
17	"""	17	"""
18	# Test with different data types	18	# Test with different data types
19	dtypes = [torch.float32, torch.float16, torch.bfloat16, torch.int32]	19	dtypes = [torch.float32, torch.float16, torch.bfloat16, torch.int32]
20	test_cases = [	20	test_cases = [
21	# (numels, axis_names, tiling_axis, no_loop_axis, split_axis, low_dims)	21	# (numels, axis_names, tiling_axis, no_loop_axis, split_axis, low_dims)
22	((1024, 512), ["h", "w"], [0, 1], [], [0, 1], []),	22	((1024, 512), ["h", "w"], [0, 1], [], [0, 1], []),
23	((2048, 2048), ["h", "w"], [0, 1], [], [0, 1], []),	23	((2048, 2048), ["h", "w"], [0, 1], [], [0, 1], []),
24	((100, 200), ["h", "w"], [0, 1], [], [0, 1], []), # Non-power-of-2 sizes	24	((100, 200), ["h", "w"], [0, 1], [], [0, 1], []), # Non-power-of-2 sizes
25	]	25	]
26		26
27	for dtype in dtypes:	27	for dtype in dtypes:
28	for numels, axis_names, tiling_axis, no_loop_axis, split_axis, low_dims in test_cases:	28	for numels, axis_names, tiling_axis, no_loop_axis, split_axis, low_dims in test_cases:
29	# Create TileGenerator instance	29	# Create TileGenerator instance
30	tile_gen = TileGenerator(	30	tile_gen = TileGenerator(
31	numels=numels,	31	numels=numels,
32	axis_names=axis_names,	32	axis_names=axis_names,
33	tiling_axis=tiling_axis,	33	tiling_axis=tiling_axis,
34	no_loop_axis=no_loop_axis,	34	no_loop_axis=no_loop_axis,
35	split_axis=split_axis,	35	split_axis=split_axis,
36	low_dims=low_dims,	36	low_dims=low_dims,
37	persistent_reduction=False,	37	persistent_reduction=False,
38	dtype=dtype,	38	dtype=dtype,
39	npu_kernel_type=NPUKernelType.SIMD	39	npu_kernel_type=NPUKernelType.SIMD
40	)	40	)
41		41
42	# Generate configs	42	# Generate configs
43	configs = tile_gen.descend_split_tiling()	43	configs = tile_gen.descend_split_tiling()
44		44
45	# Check each config for 32Byte alignment	45	# Check each config for 32Byte alignment
46	for cfg in configs:	46	for cfg in configs:
47	kwargs = cfg.kwargs	47	kwargs = cfg.kwargs
48	dtype_bytes = torch.tensor([], dtype=dtype).element_size()	48	dtype_bytes = torch.tensor([], dtype=dtype).element_size()
49	min_numel = 32 // dtype_bytes	49	min_numel = 32 // dtype_bytes
50		50
51	# Check split axis block sizes	51	# Check split axis block sizes
52	for axis in split_axis:	52	for axis in split_axis:
53	block_name = f"{axis_names[axis].upper()}BLOCK"	53	block_name = f"{axis_names[axis].upper()}BLOCK"
54	if block_name in kwargs:	54	if block_name in kwargs:
55	block_size = kwargs[block_name]	55	block_size = kwargs[block_name]
56	if block_size > min_numel:	56	if block_size > min_numel:
57	self.assertEqual(block_size % min_numel, 0,	57	self.assertEqual(block_size % min_numel, 0,
58	msg=f"Block size {block_size} for {block_name} "	58	msg=f"Block size {block_size} for {block_name} "
59	f"with dtype {dtype} is not 32Byte aligned")	59	f"with dtype {dtype} is not 32Byte aligned")
60		60
61	# Check tiling axis sub-block sizes	61	# Check tiling axis sub-block sizes
62	for axis in tiling_axis:	62	for axis in tiling_axis:
63	sub_block_name = f"{axis_names[axis].upper()}BLOCK_SUB"	63	sub_block_name = f"{axis_names[axis].upper()}BLOCK_SUB"
64	if sub_block_name in kwargs:	64	if sub_block_name in kwargs:
65	sub_block_size = kwargs[sub_block_name]	65	sub_block_size = kwargs[sub_block_name]
66	if sub_block_size > min_numel:	66	if sub_block_size > min_numel:
67	self.assertEqual(sub_block_size % min_numel, 0,	67	self.assertEqual(sub_block_size % min_numel, 0,
68	msg=f"Sub-block size {sub_block_size} for {sub_block_name} "	68	msg=f"Sub-block size {sub_block_size} for {sub_block_name} "
69	f"with dtype {dtype} is not 32Byte aligned")	69	f"with dtype {dtype} is not 32Byte aligned")
70		70
71	def test_aligned_numel_method(self):	71	def test_aligned_numel_method(self):
72	"""	72	"""
73	Test the aligned_numel_32byte method directly	73	Test the aligned_numel_32byte method directly
74	"""	74	"""
75	# Test with different data types	75	# Test with different data types
76	test_cases = [	76	test_cases = [
77	(torch.float32, 32, 32), # 32 elements = 128 bytes, already aligned	77	(torch.float32, 32, 32), # 32 elements = 128 bytes, already aligned
78	(torch.float32, 33, 40), # 33 elements = 132 bytes, should align to 40 elements (160 bytes)	78	(torch.float32, 33, 40), # 33 elements = 132 bytes, should align to 40 elements (160 bytes)
79	(torch.float16, 32, 32), # 32 elements = 64 bytes, already aligned	79	(torch.float16, 32, 32), # 32 elements = 64 bytes, already aligned
80	(torch.float16, 33, 48), # 33 elements = 66 bytes, should align to 48 elements (96 bytes)	80	(torch.float16, 33, 48), # 33 elements = 66 bytes, should align to 48 elements (96 bytes)
81	(torch.bfloat16, 32, 32), # 32 elements = 64 bytes, already aligned	81	(torch.bfloat16, 32, 32), # 32 elements = 64 bytes, already aligned
82	(torch.bfloat16, 33, 48), # 33 elements = 66 bytes, should align to 48 elements (96 bytes)	82	(torch.bfloat16, 33, 48), # 33 elements = 66 bytes, should align to 48 elements (96 bytes)
83	(torch.int32, 32, 32), # 32 elements = 128 bytes, already aligned	83	(torch.int32, 32, 32), # 32 elements = 128 bytes, already aligned
84	(torch.int32, 33, 40), # 33 elements = 132 bytes, should align to 40 elements (160 bytes)	84	(torch.int32, 33, 40), # 33 elements = 132 bytes, should align to 40 elements (160 bytes)
85	]	85	]
86		86
87	for dtype, input_numel, expected_numel in test_cases:	87	for dtype, input_numel, expected_numel in test_cases:
88	# Get dtype_bytes for the test	88	# Get dtype_bytes for the test
89	dtype_bytes = torch.tensor([], dtype=dtype).element_size()	89	dtype_bytes = torch.tensor([], dtype=dtype).element_size()
90	result = aligned_numel_32byte(input_numel, dtype_bytes)	90	result = aligned_numel_32byte(input_numel, dtype_bytes)
91	self.assertEqual(result, expected_numel,	91	self.assertEqual(result, expected_numel,
92	msg=f"aligned_numel_32byte({input_numel}) for {dtype} "	92	msg=f"aligned_numel_32byte({input_numel}) for {dtype} "
93	f"returned {result}, expected {expected_numel}")	93	f"returned {result}, expected {expected_numel}")
94		94
95		95
96	if __name__ == "__main__":	96	if __name__ == "__main__":
97	run_tests()	97	run_tests()

	@@ -1,41 +1,41 @@
1	import numpy as np	1	import numpy as np
2	import torch	2	import torch
3		3
4	import torch_npu	4	import torch_npu
5	from torch_npu.testing.testcase import TestCase, run_tests	5	from torch_npu.testing.testcase import TestCase, run_tests
6	from torch_npu.testing.common_utils import create_common_tensor	6	from torch_npu.testing.common_utils import create_common_tensor
7		7
8		8
9	class TestFastGelu(TestCase):	9	class TestFastGelu(TestCase):
10		10
11	def supported_op_exec(self, input1):	11	def supported_op_exec(self, input1):
12	attr = 1.702	12	attr = 1.702
13	attr_half = attr / 2	13	attr_half = attr / 2
14	abs_input1 = torch.abs(input1)	14	abs_input1 = torch.abs(input1)
15	numerator = input1 * torch.exp((attr_half * input1) * (input1 - abs_input1))	15	numerator = input1 * torch.exp((attr_half * input1) * (input1 - abs_input1))
16	denominator = 1.0 + torch.exp(- attr * abs_input1)	16	denominator = 1.0 + torch.exp(- attr * abs_input1)
17	output = numerator / denominator	17	output = numerator / denominator
18	return output.cpu().detach()	18	return output.cpu().detach()
19		19
20	def custom_op_exec(self, input1):	20	def custom_op_exec(self, input1):
21	output = torch_npu.fast_gelu(input1)	21	output = torch_npu.fast_gelu(input1)
22	return output.cpu().detach()	22	return output.cpu().detach()
23		23
24	def test_fast_gelu(self, device="npu"):	24	def test_fast_gelu(self, device="npu"):
25	item = [np.float32, 0, [3, 16, 32]]	25	item = [np.float32, 0, [3, 16, 32]]
26	_, npu_input = create_common_tensor(item, 0, 100)	26	_, npu_input = create_common_tensor(item, 0, 100)
27		27
28	supported_output = self.supported_op_exec(npu_input)	28	supported_output = self.supported_op_exec(npu_input)
29	custom_output = self.custom_op_exec(npu_input)	29	custom_output = self.custom_op_exec(npu_input)
30	self.assertRtolEqual(supported_output, custom_output)	30	self.assertRtolEqual(supported_output, custom_output)
31		31
32	def test_fast_gelu_input_arg(self):	32	def test_fast_gelu_input_arg(self):
33	item = [np.float32, 0, [3, 16, 32]]	33	item = [np.float32, 0, [3, 16, 32]]
34	_, npu_input = create_common_tensor(item, 0, 100)	34	_, npu_input = create_common_tensor(item, 0, 100)
35	supported_output = self.supported_op_exec(npu_input)	35	supported_output = self.supported_op_exec(npu_input)
36	custom_output = torch_npu.fast_gelu(input=npu_input)	36	custom_output = torch_npu.fast_gelu(input=npu_input)
37	self.assertRtolEqual(supported_output, custom_output)	37	self.assertRtolEqual(supported_output, custom_output)
38		38
39		39
40	if __name__ == "__main__":	40	if __name__ == "__main__":
41	run_tests()	41	run_tests()

	@@ -1,44 +1,44 @@
1	import torch	1	import torch
2		2
3	import torch_npu	3	import torch_npu
4	from torch_npu.testing.testcase import TestCase, run_tests	4	from torch_npu.testing.testcase import TestCase, run_tests
5		5
6		6
7	class TestGridAssignPositive(TestCase):	7	class TestGridAssignPositive(TestCase):
8		8
9	def supported_op_exec(self, input1, box_responsible_flags, max_overlaps, argmax_overlaps, pos_iou_thr):	9	def supported_op_exec(self, input1, box_responsible_flags, max_overlaps, argmax_overlaps, pos_iou_thr):
10	pos_inds = (max_overlaps > pos_iou_thr) & box_responsible_flags.type(torch.bool)	10	pos_inds = (max_overlaps > pos_iou_thr) & box_responsible_flags.type(torch.bool)
11	argmax_overlaps = argmax_overlaps.to(input1.dtype)	11	argmax_overlaps = argmax_overlaps.to(input1.dtype)
12	input1[pos_inds] = argmax_overlaps[pos_inds] + 1	12	input1[pos_inds] = argmax_overlaps[pos_inds] + 1
13	return input1.cpu().detach()	13	return input1.cpu().detach()
14		14
15	def custom_op_exec(self, input1, overlaps, box_responsible_flags, max_overlaps, argmax_overlaps,	15	def custom_op_exec(self, input1, overlaps, box_responsible_flags, max_overlaps, argmax_overlaps,
16	gt_max_overlaps, gt_argmax_overlaps, num_gts, pos_iou_thr, min_pos_iou, gt_max_assign_all):	16	gt_max_overlaps, gt_argmax_overlaps, num_gts, pos_iou_thr, min_pos_iou, gt_max_assign_all):
17	output = torch_npu.npu_grid_assign_positive(input1, overlaps, box_responsible_flags, max_overlaps,	17	output = torch_npu.npu_grid_assign_positive(input1, overlaps, box_responsible_flags, max_overlaps,
18	argmax_overlaps, gt_max_overlaps, gt_argmax_overlaps, num_gts,	18	argmax_overlaps, gt_max_overlaps, gt_argmax_overlaps, num_gts,
19	pos_iou_thr, min_pos_iou, gt_max_assign_all)	19	pos_iou_thr, min_pos_iou, gt_max_assign_all)
20	return output.cpu().detach()	20	return output.cpu().detach()
21		21
22	def test_npu_grid_assign_positive(self):	22	def test_npu_grid_assign_positive(self):
23	npu_input = torch.rand((4,), dtype=torch.float32).to("npu")	23	npu_input = torch.rand((4,), dtype=torch.float32).to("npu")
24	overlaps = torch.rand((2, 4), dtype=torch.float32).to("npu")	24	overlaps = torch.rand((2, 4), dtype=torch.float32).to("npu")
25	box_responsible_flags = torch.tensor([1, 1, 1, 0], dtype=torch.uint8).to("npu")	25	box_responsible_flags = torch.tensor([1, 1, 1, 0], dtype=torch.uint8).to("npu")
26	max_overlaps = torch.rand((4,), dtype=torch.float32).to("npu")	26	max_overlaps = torch.rand((4,), dtype=torch.float32).to("npu")
27	argmax_overlaps = torch.tensor([1, 0, 1, 0], dtype=torch.int32).to("npu")	27	argmax_overlaps = torch.tensor([1, 0, 1, 0], dtype=torch.int32).to("npu")
28	gt_max_overlaps = torch.rand((2,), dtype=torch.float32).to("npu")	28	gt_max_overlaps = torch.rand((2,), dtype=torch.float32).to("npu")
29	gt_argmax_overlaps = torch.tensor([1, 0], dtype=torch.int32).to("npu")	29	gt_argmax_overlaps = torch.tensor([1, 0], dtype=torch.int32).to("npu")
30	num_gts = 128	30	num_gts = 128
31	pos_iou_thr = .5	31	pos_iou_thr = .5
32	min_pos_iou = .0	32	min_pos_iou = .0
33	gt_max_assign_all = True	33	gt_max_assign_all = True
34		34
35	supported_output = self.supported_op_exec(npu_input, box_responsible_flags, max_overlaps,	35	supported_output = self.supported_op_exec(npu_input, box_responsible_flags, max_overlaps,
36	argmax_overlaps, pos_iou_thr)	36	argmax_overlaps, pos_iou_thr)
37	custom_output = self.custom_op_exec(npu_input, overlaps, box_responsible_flags, max_overlaps,	37	custom_output = self.custom_op_exec(npu_input, overlaps, box_responsible_flags, max_overlaps,
38	argmax_overlaps, gt_max_overlaps, gt_argmax_overlaps, num_gts,	38	argmax_overlaps, gt_max_overlaps, gt_argmax_overlaps, num_gts,
39	pos_iou_thr, min_pos_iou, gt_max_assign_all)	39	pos_iou_thr, min_pos_iou, gt_max_assign_all)
40	self.assertRtolEqual(supported_output, custom_output)	40	self.assertRtolEqual(supported_output, custom_output)
41		41
42		42
43	if __name__ == "__main__":	43	if __name__ == "__main__":
44	run_tests()	44	run_tests()

	@@ -1,94 +1,94 @@
1	import numpy as np	1	import numpy as np
2	import torch	2	import torch
3		3
4	import torch_npu	4	import torch_npu
5	from torch_npu.testing.testcase import TestCase, run_tests	5	from torch_npu.testing.testcase import TestCase, run_tests
6	from torch_npu.testing.common_utils import create_common_tensor	6	from torch_npu.testing.common_utils import create_common_tensor
7		7
8		8
9	class TestPSROIPooling(TestCase):	9	class TestPSROIPooling(TestCase):
10		10
11	def cal_sum(self, hstart, hend, wstart, wend, image):	11	def cal_sum(self, hstart, hend, wstart, wend, image):
12	out_sum = 0.0	12	out_sum = 0.0
13	for row in range(hstart, hend):	13	for row in range(hstart, hend):
14	for col in range(wstart, wend):	14	for col in range(wstart, wend):
15	out_sum += image[row][col]	15	out_sum += image[row][col]
16	return out_sum	16	return out_sum
17		17
18	def supported_op_exec(self, input1, rois, spatial_scale, group_size, output_dim):	18	def supported_op_exec(self, input1, rois, spatial_scale, group_size, output_dim):
19	dst_type = input1.dtype	19	dst_type = input1.dtype
20	if dst_type != torch.float32:	20	if dst_type != torch.float32:
21	input1 = input1.to(torch.float32)	21	input1 = input1.to(torch.float32)
22	rois = rois.to(torch.float32)	22	rois = rois.to(torch.float32)
23		23
24	n, channels, height, width = input1.shape	24	n, channels, height, width = input1.shape
25	tensor_height = torch.tensor([height]).npu()	25	tensor_height = torch.tensor([height]).npu()
26	tensor_width = torch.tensor([width]).npu()	26	tensor_width = torch.tensor([width]).npu()
27	tensor_zero = torch.tensor([0]).npu()	27	tensor_zero = torch.tensor([0]).npu()
28	tensor_one_tenth = torch.tensor([0.1]).npu()	28	tensor_one_tenth = torch.tensor([0.1]).npu()
29		29
30	output_size = [rois.size(0) * rois.size(2), output_dim, group_size, group_size]	30	output_size = [rois.size(0) * rois.size(2), output_dim, group_size, group_size]
31	rois = rois.transpose(2, 1)	31	rois = rois.transpose(2, 1)
32	rois = torch.reshape(rois, (rois.shape[0] * rois.shape[1], rois.shape[2]))	32	rois = torch.reshape(rois, (rois.shape[0] * rois.shape[1], rois.shape[2]))
33	output = torch.zeros(output_size).npu()	33	output = torch.zeros(output_size).npu()
34	for out_n in range(output_size[0]):	34	for out_n in range(output_size[0]):
35	for ctop in range(output_size[1]):	35	for ctop in range(output_size[1]):
36	for ph in range(output_size[2]):	36	for ph in range(output_size[2]):
37	for pw in range(output_size[3]):	37	for pw in range(output_size[3]):
38	roi_batch_ind = rois[out_n][0].to(int)	38	roi_batch_ind = rois[out_n][0].to(int)
39	roi_start_w = torch.round(rois[out_n][1]) * spatial_scale	39	roi_start_w = torch.round(rois[out_n][1]) * spatial_scale
40	roi_start_h = torch.round(rois[out_n][2]) * spatial_scale	40	roi_start_h = torch.round(rois[out_n][2]) * spatial_scale
41	roi_end_w = torch.round(rois[out_n][3] + 1.0) * spatial_scale	41	roi_end_w = torch.round(rois[out_n][3] + 1.0) * spatial_scale
42	roi_end_h = torch.round(rois[out_n][4] + 1.0) * spatial_scale	42	roi_end_h = torch.round(rois[out_n][4] + 1.0) * spatial_scale
43	# Force too small ROIs to be 1x1	43	# Force too small ROIs to be 1x1
44	roi_width = torch.max(roi_end_w - roi_start_w, tensor_one_tenth)	44	roi_width = torch.max(roi_end_w - roi_start_w, tensor_one_tenth)
45	roi_height = torch.max(roi_end_h - roi_start_h, tensor_one_tenth)	45	roi_height = torch.max(roi_end_h - roi_start_h, tensor_one_tenth)
46	# Compute w and h at bottom	46	# Compute w and h at bottom
47	bin_size_h = roi_height / group_size	47	bin_size_h = roi_height / group_size
48	bin_size_w = roi_width / group_size	48	bin_size_w = roi_width / group_size
49		49
50	# Add roi offsets and clip to input boundaries	50	# Add roi offsets and clip to input boundaries
51	hstart = (ph * bin_size_h + roi_start_h).to(int)	51	hstart = (ph * bin_size_h + roi_start_h).to(int)
52	wstart = torch.floor(pw * bin_size_w + roi_start_w).to(int)	52	wstart = torch.floor(pw * bin_size_w + roi_start_w).to(int)
53	hend = torch.ceil((ph + 1) * bin_size_h + roi_start_h).to(int)	53	hend = torch.ceil((ph + 1) * bin_size_h + roi_start_h).to(int)
54	wend = torch.ceil((pw + 1) * bin_size_w + roi_start_w).to(int)	54	wend = torch.ceil((pw + 1) * bin_size_w + roi_start_w).to(int)
55	hstart = torch.min(torch.max(hstart, tensor_zero), tensor_height).to(int)	55	hstart = torch.min(torch.max(hstart, tensor_zero), tensor_height).to(int)
56	hend = torch.min(torch.max(hend, tensor_zero), tensor_height).to(int)	56	hend = torch.min(torch.max(hend, tensor_zero), tensor_height).to(int)
57	wstart = torch.min(torch.max(wstart, tensor_zero), tensor_width).to(int)	57	wstart = torch.min(torch.max(wstart, tensor_zero), tensor_width).to(int)
58	wend = torch.min(torch.max(wend, tensor_zero), tensor_width).to(int)	58	wend = torch.min(torch.max(wend, tensor_zero), tensor_width).to(int)
59	is_empty = ((hend <= hstart) or (wend <= wstart))	59	is_empty = ((hend <= hstart) or (wend <= wstart))
60	c = (ctop * group_size + ph) * group_size + pw	60	c = (ctop * group_size + ph) * group_size + pw
61	out_sum = self.cal_sum(hstart,	61	out_sum = self.cal_sum(hstart,
62	hend,	62	hend,
63	wstart,	63	wstart,
64	wend,	64	wend,
65	image=input1[roi_batch_ind][c])	65	image=input1[roi_batch_ind][c])
66	bin_area = (hend - hstart) * (wend - wstart)	66	bin_area = (hend - hstart) * (wend - wstart)
67	if is_empty:	67	if is_empty:
68	output[out_n][ctop][ph][pw] = 0.0	68	output[out_n][ctop][ph][pw] = 0.0
69	else:	69	else:
70	output[out_n][ctop][ph][pw] = out_sum / bin_area	70	output[out_n][ctop][ph][pw] = out_sum / bin_area
71		71
72	if dst_type != torch.float32:	72	if dst_type != torch.float32:
73	output = output.to(dst_type)	73	output = output.to(dst_type)
74	return output.cpu().detach()	74	return output.cpu().detach()
75		75
76	def custom_op_exec(self, input1, rois, spatial_scale, group_size, output_dim):	76	def custom_op_exec(self, input1, rois, spatial_scale, group_size, output_dim):
77	output = torch_npu.npu_ps_roi_pooling(input1, rois, spatial_scale, group_size, output_dim)	77	output = torch_npu.npu_ps_roi_pooling(input1, rois, spatial_scale, group_size, output_dim)
78	return output.cpu().detach()	78	return output.cpu().detach()
79		79
80	def test_npu_ps_roi_pooling(self, device="npu"):	80	def test_npu_ps_roi_pooling(self, device="npu"):
81	item = [np.float32, 0, (2, 961, 127, 127)]	81	item = [np.float32, 0, (2, 961, 127, 127)]
82	_, npu_input = create_common_tensor(item, 0.1, 1)	82	_, npu_input = create_common_tensor(item, 0.1, 1)
83	rois = torch.tensor([[[0], [3], [2], [9], [9]], [[1], [1], [4], [9], [9]]], dtype=torch.float32).npu()	83	rois = torch.tensor([[[0], [3], [2], [9], [9]], [[1], [1], [4], [9], [9]]], dtype=torch.float32).npu()
84	spatial_scale = 0.25	84	spatial_scale = 0.25
85	group_size = 31	85	group_size = 31
86	output_dim = 1	86	output_dim = 1
87		87
88	supported_output = self.supported_op_exec(npu_input, rois, spatial_scale, group_size, output_dim)	88	supported_output = self.supported_op_exec(npu_input, rois, spatial_scale, group_size, output_dim)
89	custom_output = self.custom_op_exec(npu_input, rois, spatial_scale, group_size, output_dim)	89	custom_output = self.custom_op_exec(npu_input, rois, spatial_scale, group_size, output_dim)
90	self.assertRtolEqual(supported_output, custom_output)	90	self.assertRtolEqual(supported_output, custom_output)
91		91
92		92
93	if __name__ == "__main__":	93	if __name__ == "__main__":
94	run_tests()	94	run_tests()

	@@ -1,138 +1,138 @@
1	import unittest	1	import unittest
2	import torch	2	import torch
3	import numpy as np	3	import numpy as np
4		4
5	import torch_npu	5	import torch_npu
6	from torch_npu.testing.testcase import TestCase, run_tests	6	from torch_npu.testing.testcase import TestCase, run_tests
7	from torch_npu.testing.common_utils import create_common_tensor	7	from torch_npu.testing.common_utils import create_common_tensor
8		8
9		9
10	class TestCopyKernelMemoryFormat(TestCase):	10	class TestCopyKernelMemoryFormat(TestCase):
11	def test_h2d_copy_contiguous_tensor(self):	11	def test_h2d_copy_contiguous_tensor(self):
12	dtype_list = [np.float16, np.float32, np.int32, np.int64]	12	dtype_list = [np.float16, np.float32, np.int32, np.int64]
13	shape_list = [[10, 20], [32, 64, 128], [2, 3, 4, 5]]	13	shape_list = [[10, 20], [32, 64, 128], [2, 3, 4, 5]]
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()
21	self.assertRtolEqual(npu_input_copy.cpu().numpy(), cpu_input.numpy())	21	self.assertRtolEqual(npu_input_copy.cpu().numpy(), cpu_input.numpy())
22		22
23	def test_h2d_copy_non_contiguous_tensor(self):	23	def test_h2d_copy_non_contiguous_tensor(self):
24	dtype_list = [np.float16, np.float32]	24	dtype_list = [np.float16, np.float32]
25	shape_list = [[32, 64], [16, 32, 64]]	25	shape_list = [[32, 64], [16, 32, 64]]
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)
33	npu_transposed = cpu_transposed.npu()	33	npu_transposed = cpu_transposed.npu()
34	npu_contiguous = npu_transposed.contiguous()	34	npu_contiguous = npu_transposed.contiguous()
35	self.assertRtolEqual(npu_contiguous.cpu().numpy(), cpu_transposed.contiguous().numpy())	35	self.assertRtolEqual(npu_contiguous.cpu().numpy(), cpu_transposed.contiguous().numpy())
36		36
37	def test_d2h_copy_contiguous_tensor(self):	37	def test_d2h_copy_contiguous_tensor(self):
38	dtype_list = [np.float16, np.float32, np.int32, np.int64]	38	dtype_list = [np.float16, np.float32, np.int32, np.int64]
39	shape_list = [[10, 20], [32, 64, 128], [2, 3, 4, 5]]	39	shape_list = [[10, 20], [32, 64, 128], [2, 3, 4, 5]]
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()
47	self.assertRtolEqual(cpu_output.numpy(), cpu_input.numpy())	47	self.assertRtolEqual(cpu_output.numpy(), cpu_input.numpy())
48		48
49	def test_d2h_copy_non_contiguous_tensor(self):	49	def test_d2h_copy_non_contiguous_tensor(self):
50	dtype_list = [np.float16, np.float32]	50	dtype_list = [np.float16, np.float32]
51	shape_list = [[32, 64], [16, 32, 64]]	51	shape_list = [[32, 64], [16, 32, 64]]
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)
59	cpu_output = npu_transposed.cpu()	59	cpu_output = npu_transposed.cpu()
60	self.assertRtolEqual(cpu_output.numpy(), cpu_input.transpose(-1, -2).contiguous().numpy())	60	self.assertRtolEqual(cpu_output.numpy(), cpu_input.transpose(-1, -2).contiguous().numpy())
61		61
62	def test_h2d_copy_different_dtype(self):	62	def test_h2d_copy_different_dtype(self):
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
77	def test_d2h_copy_different_dtype(self):	77	def test_d2h_copy_different_dtype(self):
78	dtype_pairs = [	78	dtype_pairs = [
79	(np.float16, torch.float32),	79	(np.float16, torch.float32),
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
101	def test_d2h_copy_slice_tensor(self):	101	def test_d2h_copy_slice_tensor(self):
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())
109		109
110	def test_d2h_copy_broadcast_tensor(self):	110	def test_d2h_copy_broadcast_tensor(self):
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())
118		118
119	def test_h2d_copy_permute_tensor(self):	119	def test_h2d_copy_permute_tensor(self):
120	shape = [32, 64, 128]	120	shape = [32, 64, 128]
121	cpu_input = torch.randn(shape)	121	cpu_input = torch.randn(shape)
122	cpu_permuted = cpu_input.permute(2, 0, 1)	122	cpu_permuted = cpu_input.permute(2, 0, 1)
123	npu_permuted = cpu_permuted.npu()	123	npu_permuted = cpu_permuted.npu()
124	npu_contiguous = npu_permuted.contiguous()	124	npu_contiguous = npu_permuted.contiguous()
125	self.assertRtolEqual(npu_contiguous.cpu().numpy(), cpu_permuted.contiguous().numpy())	125	self.assertRtolEqual(npu_contiguous.cpu().numpy(), cpu_permuted.contiguous().numpy())
126		126
127	def test_d2h_copy_permute_tensor(self):	127	def test_d2h_copy_permute_tensor(self):
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
136		136
137	if __name__ == "__main__":	137	if __name__ == "__main__":
138	run_tests()	138	run_tests()

	@@ -1,19 +1,19 @@
1	import unittest	1	import unittest
2		2
3	import os	3	import os
4	import torch	4	import torch
5	from torch.testing._internal.common_utils import TestCase, run_tests	5	from torch.testing._internal.common_utils import TestCase, run_tests
6	from torch_npu.npu.utils import get_cann_version	6	from torch_npu.npu.utils import get_cann_version
7		7
8		8
9	class TestPinMemoryHostRegister(TestCase):	9	class TestPinMemoryHostRegister(TestCase):
10		10
11	@unittest.skipUnless(get_cann_version("RUNTIME") >= "8.5.0" and get_cann_version(module="DRIVER") >= "25.5.0", "This feature is not supported in older versions.")	11	@unittest.skipUnless(get_cann_version("RUNTIME") >= "8.5.0" and get_cann_version(module="DRIVER") >= "25.5.0", "This feature is not supported in older versions.")
12	def test_pin_memory_host_register(self):	12	def test_pin_memory_host_register(self):
13	os.environ["PYTORCH_NPU_ALLOC_CONF"] = "pinned_mem_register:True"	13	os.environ["PYTORCH_NPU_ALLOC_CONF"] = "pinned_mem_register:True"
14	cpu_tensor = torch.ones([2, 3])	14	cpu_tensor = torch.ones([2, 3])
15	pin_tensor = cpu_tensor.pin_memory()	15	pin_tensor = cpu_tensor.pin_memory()
16	self.assertTrue(pin_tensor.is_pinned())	16	self.assertTrue(pin_tensor.is_pinned())
17		17
18	if __name__ == '__main__':	18	if __name__ == '__main__':
19	run_tests()	19	run_tests()

	@@ -1,44 +1,44 @@
1	#include "torch_npu/csrc/afd/Init.h"	1	#include "torch_npu/csrc/afd/Init.h"
2	#include <torch/csrc/python_headers.h>	2	#include <torch/csrc/python_headers.h>
3	#include <torch/csrc/Exceptions.h>	3	#include <torch/csrc/Exceptions.h>
4	#include <torch/csrc/utils/object_ptr.h>	4	#include <torch/csrc/utils/object_ptr.h>
5	#include <torch/csrc/utils/pybind.h>	5	#include <torch/csrc/utils/pybind.h>
6	#include "torch_npu/csrc/afd/ScheduleContext.h"	6	#include "torch_npu/csrc/afd/ScheduleContext.h"
7		7
8	namespace torch_npu {	8	namespace torch_npu {
9	namespace afd {	9	namespace afd {
10		10
11	PyObject afd_init(PyObject _unused, PyObject * noargs)	11	PyObject afd_init(PyObject _unused, PyObject * noargs)
12	{	12	{
13	auto torch_npu_C_module = THPObjectPtr(PyImport_ImportModule("torch_npu._C"));	13	auto torch_npu_C_module = THPObjectPtr(PyImport_ImportModule("torch_npu._C"));
14	if (!torch_npu_C_module) {	14	if (!torch_npu_C_module) {
15	throw python_error();	15	throw python_error();
16	}	16	}
17	auto torch_npu_C_m = py::handle(torch_npu_C_module).cast<py::module>();	17	auto torch_npu_C_m = py::handle(torch_npu_C_module).cast<py::module>();
18		18
19	auto m = torch_npu_C_m.def_submodule("_afd", "Attention-FFN Disaggregation");	19	auto m = torch_npu_C_m.def_submodule("_afd", "Attention-FFN Disaggregation");
20	auto module = py::handle(m).cast<py::module>();	20	auto module = py::handle(m).cast<py::module>();
21		21
22	py::class_<ScheduleContextHolder>(module, "ScheduleContextHolder")	22	py::class_<ScheduleContextHolder>(module, "ScheduleContextHolder")
23	.def(py::init<int32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint64_t, uint64_t,	23	.def(py::init<int32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint64_t, uint64_t,
24	uint64_t, uint64_t>())	24	uint64_t, uint64_t>())
25	.def("init", &ScheduleContextHolder::Init)	25	.def("init", &ScheduleContextHolder::Init)
26	.def("get_context_tensor", &ScheduleContextHolder::GetContextTensor)	26	.def("get_context_tensor", &ScheduleContextHolder::GetContextTensor)
27	.def("stop_schedule", &ScheduleContextHolder::StopSchedule)	27	.def("stop_schedule", &ScheduleContextHolder::StopSchedule)
28	.def("get_schedule_context_info", &ScheduleContextHolder::GetScheduleContextInfo);	28	.def("get_schedule_context_info", &ScheduleContextHolder::GetScheduleContextInfo);
29	Py_RETURN_TRUE;	29	Py_RETURN_TRUE;
30	}	30	}
31		31
32	// methods on torch._C	32	// methods on torch._C
33	PyMethodDef methods[] = {	33	PyMethodDef methods[] = {
34	{"_afd_init", afd_init, METH_NOARGS, nullptr},	34	{"_afd_init", afd_init, METH_NOARGS, nullptr},
35	{nullptr, nullptr, 0, nullptr}	35	{nullptr, nullptr, 0, nullptr}
36	};	36	};
37		37
38	PyMethodDef *python_functions()	38	PyMethodDef *python_functions()
39	{	39	{
40	return methods;	40	return methods;
41	}	41	}
42		42
43	} // namespace afd	43	} // namespace afd
44	} // namespace torch_npu	44	} // namespace torch_npu

	@@ -1,311 +1,311 @@
1	#include <filesystem>	1	#include <filesystem>
2	#include <fstream>	2	#include <fstream>
3	#include <string>	3	#include <string>
4		4
5	#include <torch/csrc/distributed/c10d/Utils.hpp>	5	#include <torch/csrc/distributed/c10d/Utils.hpp>
6		6
7	#include "torch_npu/csrc/core/npu/interface/HcclInterface.h"	7	#include "torch_npu/csrc/core/npu/interface/HcclInterface.h"
8	#include "torch_npu/csrc/distributed/HCCLUtils.hpp"	8	#include "torch_npu/csrc/distributed/HCCLUtils.hpp"
9		9
10		10
11	namespace c10d_npu {	11	namespace c10d_npu {
12	bool isFileExists(const std::string& path)	12	bool isFileExists(const std::string& path)
13	{	13	{
14	std::filesystem::path filePath(path);	14	std::filesystem::path filePath(path);
15		15
16	if (!filePath.is_absolute()) {	16	if (!filePath.is_absolute()) {
17	TORCH_CHECK(false, "Path is not absolute.", DIST_ERROR(ErrCode::UNAVAIL))	17	TORCH_CHECK(false, "Path is not absolute.", DIST_ERROR(ErrCode::UNAVAIL))
18	return false;	18	return false;
19	}	19	}
20		20
21	if (std::filesystem::exists(filePath) && std::filesystem::is_regular_file(filePath)) {	21	if (std::filesystem::exists(filePath) && std::filesystem::is_regular_file(filePath)) {
22	return true;	22	return true;
23	} else {	23	} else {
24	return false;	24	return false;
25	}	25	}
26	}	26	}
27		27
28	bool checkFilePathReadable(const std::string& file)	28	bool checkFilePathReadable(const std::string& file)
29	{	29	{
30	std::filesystem::path filePath(file);	30	std::filesystem::path filePath(file);
31		31
32	if (!std::filesystem::exists(filePath)) {	32	if (!std::filesystem::exists(filePath)) {
33	return false;	33	return false;
34	}	34	}
35		35
36	if (std::filesystem::is_symlink(filePath)) {	36	if (std::filesystem::is_symlink(filePath)) {
37	return false;	37	return false;
38	}	38	}
39		39
40	if (!std::filesystem::is_regular_file(filePath)) {	40	if (!std::filesystem::is_regular_file(filePath)) {
41	return false;	41	return false;
42	}	42	}
43		43
44	std::filesystem::perms perms = std::filesystem::status(filePath).permissions();	44	std::filesystem::perms perms = std::filesystem::status(filePath).permissions();
45	if ((perms & std::filesystem::perms::owner_read) == std::filesystem::perms::owner_read) {	45	if ((perms & std::filesystem::perms::owner_read) == std::filesystem::perms::owner_read) {
46	return true;	46	return true;
47	}	47	}
48	return false;	48	return false;
49	}	49	}
50		50
51	// HCCL DataType mapping	51	// HCCL DataType mapping
52	std::map<at::ScalarType, HcclDataType> kScalarTypeToHcclDataType = {	52	std::map<at::ScalarType, HcclDataType> kScalarTypeToHcclDataType = {
53	{at::kByte, HCCL_DATA_TYPE_UINT8},	53	{at::kByte, HCCL_DATA_TYPE_UINT8},
54	{at::kChar, HCCL_DATA_TYPE_INT8},	54	{at::kChar, HCCL_DATA_TYPE_INT8},
55	{at::kShort, HCCL_DATA_TYPE_INT16},	55	{at::kShort, HCCL_DATA_TYPE_INT16},
56	{at::kInt, HCCL_DATA_TYPE_INT32},	56	{at::kInt, HCCL_DATA_TYPE_INT32},
57	{at::kLong, HCCL_DATA_TYPE_INT64},	57	{at::kLong, HCCL_DATA_TYPE_INT64},
58	{at::kHalf, HCCL_DATA_TYPE_FP16},	58	{at::kHalf, HCCL_DATA_TYPE_FP16},
59	{at::kFloat, HCCL_DATA_TYPE_FP32},	59	{at::kFloat, HCCL_DATA_TYPE_FP32},
60	{at::ScalarType::UInt16, HCCL_DATA_TYPE_UINT16},	60	{at::ScalarType::UInt16, HCCL_DATA_TYPE_UINT16},
61	{at::ScalarType::UInt32, HCCL_DATA_TYPE_UINT32},	61	{at::ScalarType::UInt32, HCCL_DATA_TYPE_UINT32},
62	{at::ScalarType::UInt64, HCCL_DATA_TYPE_UINT64},	62	{at::ScalarType::UInt64, HCCL_DATA_TYPE_UINT64},
63	{at::kDouble, HCCL_DATA_TYPE_FP64},	63	{at::kDouble, HCCL_DATA_TYPE_FP64},
64	{at::kBool, HCCL_DATA_TYPE_UINT8},	64	{at::kBool, HCCL_DATA_TYPE_UINT8},
65	{at::kBFloat16, HCCL_DATA_TYPE_BFP16},	65	{at::kBFloat16, HCCL_DATA_TYPE_BFP16},
66	{at::ScalarType::Float8_e4m3fn, HCCL_DATA_TYPE_FP8E4M3},	66	{at::ScalarType::Float8_e4m3fn, HCCL_DATA_TYPE_FP8E4M3},
67	{at::ScalarType::Float8_e5m2, HCCL_DATA_TYPE_FP8E5M2},	67	{at::ScalarType::Float8_e5m2, HCCL_DATA_TYPE_FP8E5M2},
68	};	68	};
69		69
70	std::map<HcclDataType, std::string> kHcclDataTypeToStringMap = {	70	std::map<HcclDataType, std::string> kHcclDataTypeToStringMap = {
71	{HCCL_DATA_TYPE_UINT8, "at::kByte/at::kBool"},	71	{HCCL_DATA_TYPE_UINT8, "at::kByte/at::kBool"},
72	{HCCL_DATA_TYPE_INT8, "at::kChar"},	72	{HCCL_DATA_TYPE_INT8, "at::kChar"},
73	{HCCL_DATA_TYPE_INT16, "at::kShort"},	73	{HCCL_DATA_TYPE_INT16, "at::kShort"},
74	{HCCL_DATA_TYPE_INT32, "at::kInt"},	74	{HCCL_DATA_TYPE_INT32, "at::kInt"},
75	{HCCL_DATA_TYPE_UINT16, "at::ScalarType::UInt16"},	75	{HCCL_DATA_TYPE_UINT16, "at::ScalarType::UInt16"},
76	{HCCL_DATA_TYPE_UINT32, "at::ScalarType::UInt32"},	76	{HCCL_DATA_TYPE_UINT32, "at::ScalarType::UInt32"},
77	{HCCL_DATA_TYPE_UINT64, "at::ScalarType::UInt64"},	77	{HCCL_DATA_TYPE_UINT64, "at::ScalarType::UInt64"},
78	{HCCL_DATA_TYPE_INT64, "at::kLong"},	78	{HCCL_DATA_TYPE_INT64, "at::kLong"},
79	{HCCL_DATA_TYPE_FP16, "at::kHalf"},	79	{HCCL_DATA_TYPE_FP16, "at::kHalf"},
80	{HCCL_DATA_TYPE_FP32, "at::kFloat"},	80	{HCCL_DATA_TYPE_FP32, "at::kFloat"},
81	{HCCL_DATA_TYPE_FP64, "at::kDouble"},	81	{HCCL_DATA_TYPE_FP64, "at::kDouble"},
82	{HCCL_DATA_TYPE_BFP16, "at::kBFloat16"},	82	{HCCL_DATA_TYPE_BFP16, "at::kBFloat16"},
83	{HCCL_DATA_TYPE_FP8E4M3, "at::ScalarType::Float8_e4m3fn"},	83	{HCCL_DATA_TYPE_FP8E4M3, "at::ScalarType::Float8_e4m3fn"},
84	{HCCL_DATA_TYPE_FP8E5M2, "at::ScalarType::Float8_e5m2"},	84	{HCCL_DATA_TYPE_FP8E5M2, "at::ScalarType::Float8_e5m2"},
85	};	85	};
86		86
87	// Helper function that gets the data type and issues error if not supported	87	// Helper function that gets the data type and issues error if not supported
88	HcclDataType getHcclDataType(at::ScalarType type)	88	HcclDataType getHcclDataType(at::ScalarType type)
89	{	89	{
90	try {	90	try {
91	return kScalarTypeToHcclDataType.at(type);	91	return kScalarTypeToHcclDataType.at(type);
92	} catch (std::out_of_range& e) {	92	} catch (std::out_of_range& e) {
93	throw std::runtime_error("Unsupported data type for HCCL process group" + DIST_ERROR(ErrCode::NOT_SUPPORT));	93	throw std::runtime_error("Unsupported data type for HCCL process group" + DIST_ERROR(ErrCode::NOT_SUPPORT));
94	}	94	}
95	}	95	}
96		96
97	std::string getHcclDataTypeSerialString(HcclDataType type)	97	std::string getHcclDataTypeSerialString(HcclDataType type)
98	{	98	{
99	const auto& iter = kHcclDataTypeToStringMap.find(type);	99	const auto& iter = kHcclDataTypeToStringMap.find(type);
100	if (iter != kHcclDataTypeToStringMap.cend()) {	100	if (iter != kHcclDataTypeToStringMap.cend()) {
101	return iter->second;	101	return iter->second;
102	} else {	102	} else {
103	TORCH_NPU_WARN_ONCE("Can not serialize undefined hccl data type.");	103	TORCH_NPU_WARN_ONCE("Can not serialize undefined hccl data type.");
104	return "";	104	return "";
105	}	105	}
106	}	106	}
107		107
108	bool isSupportHcclCommName()	108	bool isSupportHcclCommName()
109	{	109	{
110	return at_npu::hccl::isHcclFeatureSupported(HcclCommConfigCapability::HCCL_COMM_CONFIG_COMM_NAME);	110	return at_npu::hccl::isHcclFeatureSupported(HcclCommConfigCapability::HCCL_COMM_CONFIG_COMM_NAME);
111	}	111	}
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();
119	}	119	}
120		120
121	std::shared_ptr<HCCLComm> HCCLComm::create(	121	std::shared_ptr<HCCLComm> HCCLComm::create(
122	int numRanks,	122	int numRanks,
123	int rank,	123	int rank,
124	HcclRootInfo& rootInfo)	124	HcclRootInfo& rootInfo)
125	{	125	{
126	auto comm = std::make_shared<HCCLComm>();	126	auto comm = std::make_shared<HCCLComm>();
127	HCCL_CHECK_ERROR(hcclCommInitRootInfo(numRanks, &rootInfo, rank, &(comm->hcclComm_)));	127	HCCL_CHECK_ERROR(hcclCommInitRootInfo(numRanks, &rootInfo, rank, &(comm->hcclComm_)));
128	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {comm->destroyHcclComm();},	128	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {comm->destroyHcclComm();},
129	c10_npu::ReleasePriority::PriorityMiddle);	129	c10_npu::ReleasePriority::PriorityMiddle);
130	return comm;	130	return comm;
131	}	131	}
132		132
133	std::shared_ptr<HCCLComm> HCCLComm::create_config(	133	std::shared_ptr<HCCLComm> HCCLComm::create_config(
134	int numRanks,	134	int numRanks,
135	int rank,	135	int rank,
136	HcclRootInfo& rootInfo,	136	HcclRootInfo& rootInfo,
137	HcclCommConfig* config)	137	HcclCommConfig* config)
138	{	138	{
139	auto comm = std::make_shared<HCCLComm>();	139	auto comm = std::make_shared<HCCLComm>();
140	HCCL_CHECK_ERROR(hcclCommInitRootInfoConfig(numRanks, &rootInfo, rank, config, &(comm->hcclComm_)));	140	HCCL_CHECK_ERROR(hcclCommInitRootInfoConfig(numRanks, &rootInfo, rank, config, &(comm->hcclComm_)));
141	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {comm->destroyHcclComm();},	141	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {comm->destroyHcclComm();},
142	c10_npu::ReleasePriority::PriorityMiddle);	142	c10_npu::ReleasePriority::PriorityMiddle);
143	return comm;	143	return comm;
144	}	144	}
145		145
146	std::shared_ptr<HCCLComm> HCCLComm::createGlobalHcclComm(	146	std::shared_ptr<HCCLComm> HCCLComm::createGlobalHcclComm(
147	const char *clusterInfo,	147	const char *clusterInfo,
148	uint32_t rank,	148	uint32_t rank,
149	HcclCommConfig* config)	149	HcclCommConfig* config)
150	{	150	{
151	auto comm = std::make_shared<HCCLComm>();	151	auto comm = std::make_shared<HCCLComm>();
152	if (hcclCommInitClusterInfoConfig(clusterInfo, rank, config, &(comm->hcclComm_)) != HCCL_SUCCESS) {	152	if (hcclCommInitClusterInfoConfig(clusterInfo, rank, config, &(comm->hcclComm_)) != HCCL_SUCCESS) {
153	return nullptr;	153	return nullptr;
154	}	154	}
155	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {comm->destroyHcclComm();},	155	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {comm->destroyHcclComm();},
156	c10_npu::ReleasePriority::PriorityMiddle);	156	c10_npu::ReleasePriority::PriorityMiddle);
157	return comm;	157	return comm;
158	}	158	}
159		159
160	std::shared_ptr<HCCLComm> HCCLComm::createSubHcclComm(	160	std::shared_ptr<HCCLComm> HCCLComm::createSubHcclComm(
161	std::shared_ptr<HCCLComm> comm,	161	std::shared_ptr<HCCLComm> comm,
162	uint32_t rankNum,	162	uint32_t rankNum,
163	uint32_t *rankIds,	163	uint32_t *rankIds,
164	uint64_t subCommId,	164	uint64_t subCommId,
165	uint32_t subCommRankId,	165	uint32_t subCommRankId,
166	HcclCommConfig* config)	166	HcclCommConfig* config)
167	{	167	{
168	auto subComm = std::make_shared<HCCLComm>();	168	auto subComm = std::make_shared<HCCLComm>();
169	if (hcclCreateSubCommConfig(&(comm->hcclComm_), rankNum, rankIds, subCommId, subCommRankId,	169	if (hcclCreateSubCommConfig(&(comm->hcclComm_), rankNum, rankIds, subCommId, subCommRankId,
170	config, &(subComm->hcclComm_)) != HCCL_SUCCESS) {	170	config, &(subComm->hcclComm_)) != HCCL_SUCCESS) {
171	return nullptr;	171	return nullptr;
172	}	172	}
173	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {subComm->destroyHcclComm();},	173	c10_npu::NpuSysCtrl::GetInstance().RegisterReleaseFn([=]() ->void {subComm->destroyHcclComm();},
174	c10_npu::ReleasePriority::PriorityMiddle);	174	c10_npu::ReleasePriority::PriorityMiddle);
175	return subComm;	175	return subComm;
176	}	176	}
177		177
178	// Move constructable	178	// Move constructable
179	HCCLComm::HCCLComm(HCCLComm&& other)	179	HCCLComm::HCCLComm(HCCLComm&& other)
180	{	180	{
181	std::swap(hcclComm_, other.hcclComm_);	181	std::swap(hcclComm_, other.hcclComm_);
182	std::swap(hcclAsyncErr_, other.hcclAsyncErr_);	182	std::swap(hcclAsyncErr_, other.hcclAsyncErr_);
183	std::swap(hcclCommType, other.hcclCommType);	183	std::swap(hcclCommType, other.hcclCommType);
184	std::swap(p2pPeer, other.p2pPeer);	184	std::swap(p2pPeer, other.p2pPeer);
185	}	185	}
186		186
187	// Move assignable	187	// Move assignable
188	HCCLComm& HCCLComm::operator=(HCCLComm&& other)	188	HCCLComm& HCCLComm::operator=(HCCLComm&& other)
189	{	189	{
190	std::swap(hcclComm_, other.hcclComm_);	190	std::swap(hcclComm_, other.hcclComm_);
191	std::swap(hcclAsyncErr_, other.hcclAsyncErr_);	191	std::swap(hcclAsyncErr_, other.hcclAsyncErr_);
192	std::swap(hcclCommType, other.hcclCommType);	192	std::swap(hcclCommType, other.hcclCommType);
193	std::swap(p2pPeer, other.p2pPeer);	193	std::swap(p2pPeer, other.p2pPeer);
194	return *this;	194	return *this;
195	}	195	}
196		196
197	void HCCLComm::destroyHcclComm()	197	void HCCLComm::destroyHcclComm()
198	{	198	{
199	std::unique_lock<std::mutex> lock(mutex_);	199	std::unique_lock<std::mutex> lock(mutex_);
200	if (hcclComm_) {	200	if (hcclComm_) {
201	hcclCommDestroy(hcclComm_);	201	hcclCommDestroy(hcclComm_);
202	hcclComm_ = nullptr;	202	hcclComm_ = nullptr;
203	}	203	}
204	}	204	}
205		205
206	HcclResult HCCLComm::checkForHcclError()	206	HcclResult HCCLComm::checkForHcclError()
207	{	207	{
208	std::unique_lock<std::mutex> lock(mutex_);	208	std::unique_lock<std::mutex> lock(mutex_);
209	#ifdef ENABLE_HCCL_ERROR_CHECKING	209	#ifdef ENABLE_HCCL_ERROR_CHECKING
210	if (hcclAsyncErr_ != HCCL_SUCCESS) {	210	if (hcclAsyncErr_ != HCCL_SUCCESS) {
211	return hcclAsyncErr_;	211	return hcclAsyncErr_;
212	}	212	}
213	if (hcclComm_ != nullptr) {	213	if (hcclComm_ != nullptr) {
214	HcclResult result = hcclGetCommAsyncError(hcclComm_, &hcclAsyncErr_);	214	HcclResult result = hcclGetCommAsyncError(hcclComm_, &hcclAsyncErr_);
215	if (result != HCCL_SUCCESS) {	215	if (result != HCCL_SUCCESS) {
216	std::string temp_str = std::string("Failed to get HCCL error code: ") + std::to_string(result);	216	std::string temp_str = std::string("Failed to get HCCL error code: ") + std::to_string(result);
217	const char* errmsg = temp_str.c_str();	217	const char* errmsg = temp_str.c_str();
218	ASCEND_LOGE("%s", errmsg);	218	ASCEND_LOGE("%s", errmsg);
219	LOG(ERROR) << c10::str(errmsg);	219	LOG(ERROR) << c10::str(errmsg);
220	return result; // return this error result instead of hcclAsyncErr_	220	return result; // return this error result instead of hcclAsyncErr_
221	}	221	}
222	}	222	}
223	return hcclAsyncErr_;	223	return hcclAsyncErr_;
224	#else	224	#else
225	// Always return success, if error checks are disabled.	225	// Always return success, if error checks are disabled.
226	return HCCL_SUCCESS;	226	return HCCL_SUCCESS;
227	#endif	227	#endif
228	}	228	}
229		229
230	void DebugInfoWriter::write(const std::string &hcclTrace)	230	void DebugInfoWriter::write(const std::string &hcclTrace)
231	{	231	{
232	// Open a file for writing. The ios::binary flag is used to write data as	232	// Open a file for writing. The ios::binary flag is used to write data as
233	// binary.	233	// binary.
234	std::ofstream file(filename_, std::ios::binary);	234	std::ofstream file(filename_, std::ios::binary);
235		235
236	// Check if the file was opened successfully.	236	// Check if the file was opened successfully.
237	if (!file.is_open()) {	237	if (!file.is_open()) {
238	LOG(ERROR) << "Error opening file for writing HCCLPG debug info: "	238	LOG(ERROR) << "Error opening file for writing HCCLPG debug info: "
239	<< filename_;	239	<< filename_;
240	return;	240	return;
241	}	241	}
242		242
243	file.write(hcclTrace.data(), hcclTrace.size());	243	file.write(hcclTrace.data(), hcclTrace.size());
244	LOG(INFO) << "Finished writing HCCLPG debug info to " << filename_;	244	LOG(INFO) << "Finished writing HCCLPG debug info to " << filename_;
245	}	245	}
246		246
247	DebugInfoWriter &DebugInfoWriter::getWriter(int rank)	247	DebugInfoWriter &DebugInfoWriter::getWriter(int rank)
248	{	248	{
249	if (writer_ == nullptr) {	249	if (writer_ == nullptr) {
250	std::string fileNamePrefix = c10d::getCvarString(	250	std::string fileNamePrefix = c10d::getCvarString(
251	{"TORCH_HCCL_DEBUG_INFO_TEMP_FILE"}, "/tmp/hccl_trace_rank_");	251	{"TORCH_HCCL_DEBUG_INFO_TEMP_FILE"}, "/tmp/hccl_trace_rank_");
252	// Using std::unique_ptr here to auto-delete the writer object	252	// Using std::unique_ptr here to auto-delete the writer object
253	// when the pointer itself is destroyed.	253	// when the pointer itself is destroyed.
254	std::unique_ptr<DebugInfoWriter> writerPtr(	254	std::unique_ptr<DebugInfoWriter> writerPtr(
255	new DebugInfoWriter(fileNamePrefix, rank));	255	new DebugInfoWriter(fileNamePrefix, rank));
256	DebugInfoWriter::registerWriter(std::move(writerPtr));	256	DebugInfoWriter::registerWriter(std::move(writerPtr));
257	}	257	}
258	return *writer_;	258	return *writer_;
259	}	259	}
260		260
261	void DebugInfoWriter::registerWriter(std::unique_ptr<DebugInfoWriter> writer)	261	void DebugInfoWriter::registerWriter(std::unique_ptr<DebugInfoWriter> writer)
262	{	262	{
263	TORCH_CHECK_WITH(	263	TORCH_CHECK_WITH(
264	DistBackendError,	264	DistBackendError,
265	!hasWriterRegistered_.load(),	265	!hasWriterRegistered_.load(),
266	"debugInfoWriter already registered");	266	"debugInfoWriter already registered");
267	hasWriterRegistered_.store(true);	267	hasWriterRegistered_.store(true);
268	writer_ = std::move(writer);	268	writer_ = std::move(writer);
269	}	269	}
270		270
271	std::unique_ptr<DebugInfoWriter> DebugInfoWriter::writer_ = nullptr;	271	std::unique_ptr<DebugInfoWriter> DebugInfoWriter::writer_ = nullptr;
272	std::atomic<bool> DebugInfoWriter::hasWriterRegistered_(false);	272	std::atomic<bool> DebugInfoWriter::hasWriterRegistered_(false);
273		273
274	struct HcclBufferNameKey {	274	struct HcclBufferNameKey {
275	c10::DeviceIndex device_index;	275	c10::DeviceIndex device_index;
276	std::string name;	276	std::string name;
277	bool operator<(const HcclBufferNameKey& other) const	277	bool operator<(const HcclBufferNameKey& other) const
278	{	278	{
279	if (device_index != other.device_index) {	279	if (device_index != other.device_index) {
280	return device_index < other.device_index;	280	return device_index < other.device_index;
281	}	281	}
282	return name < other.name; // sort by name if device_index is the same	282	return name < other.name; // sort by name if device_index is the same
283	}	283	}
284	};	284	};
285		285
286	struct HcclBufferNameStreamMap {	286	struct HcclBufferNameStreamMap {
287	std::map<HcclBufferNameKey, c10_npu::NPUStream> map;	287	std::map<HcclBufferNameKey, c10_npu::NPUStream> map;
288	std::mutex mutex;	288	std::mutex mutex;
289	} g_BufferNameStreamMap = {};	289	} g_BufferNameStreamMap = {};
290		290
291	c10::optional<c10_npu::NPUStream> getHcclStreamByBufferName(const std::string &name, c10::DeviceIndex device_index)	291	c10::optional<c10_npu::NPUStream> getHcclStreamByBufferName(const std::string &name, c10::DeviceIndex device_index)
292	{	292	{
293	std::unique_lock<std::mutex> lock(g_BufferNameStreamMap.mutex);	293	std::unique_lock<std::mutex> lock(g_BufferNameStreamMap.mutex);
294	auto &map = g_BufferNameStreamMap.map;	294	auto &map = g_BufferNameStreamMap.map;
295	auto it = map.find({device_index, name});	295	auto it = map.find({device_index, name});
296	if (it == map.end()) {	296	if (it == map.end()) {
297	return {};	297	return {};
298	}	298	}
299	return it->second;	299	return it->second;
300	}	300	}
301		301
302	bool setHcclStreamByBufferName(const std::string &name, c10::DeviceIndex device_index, c10_npu::NPUStream steam)	302	bool setHcclStreamByBufferName(const std::string &name, c10::DeviceIndex device_index, c10_npu::NPUStream steam)
303	{	303	{
304	HcclBufferNameKey key = {device_index, name};	304	HcclBufferNameKey key = {device_index, name};
305	std::unique_lock<std::mutex> lock(g_BufferNameStreamMap.mutex);	305	std::unique_lock<std::mutex> lock(g_BufferNameStreamMap.mutex);
306	auto &map = g_BufferNameStreamMap.map;	306	auto &map = g_BufferNameStreamMap.map;
307	auto pair = map.insert({key, steam});	307	auto pair = map.insert({key, steam});
308	return pair.second;	308	return pair.second;
309	}	309	}
310		310
311	} // namespace c10d_npu	311	} // namespace c10d_npu

	@@ -1,397 +1,396 @@
1	cmake_minimum_required(VERSION 3.18 FATAL_ERROR)	1	cmake_minimum_required(VERSION 3.18 FATAL_ERROR)
2		2
3	find_program(CCACHE ccache)	3	find_program(CCACHE ccache)
4	if(${CCACHE} STREQUAL "CCACHE-NOTFOUND")	4	if(${CCACHE} STREQUAL "CCACHE-NOTFOUND")
5	message(STATUS "Compile without ccache")	5	message(STATUS "Compile without ccache")
6	else()	6	else()
7	set(CMAKE_C_COMPILER_LAUNCHER ${CCACHE} CACHE PATH "cache Compiler")	7	set(CMAKE_C_COMPILER_LAUNCHER ${CCACHE} CACHE PATH "cache Compiler")
8	set(CMAKE_CXX_COMPILER_LAUNCHER ${CCACHE} CACHE PATH "cache Compiler")	8	set(CMAKE_CXX_COMPILER_LAUNCHER ${CCACHE} CACHE PATH "cache Compiler")
9	message(STATUS "CMAKE_C_COMPILER_LAUNCHER:${CMAKE_C_COMPILER_LAUNCHER}")	9	message(STATUS "CMAKE_C_COMPILER_LAUNCHER:${CMAKE_C_COMPILER_LAUNCHER}")
10	message(STATUS "CMAKE_CXX_COMPILER_LAUNCHER:${CMAKE_CXX_COMPILER_LAUNCHER}")	10	message(STATUS "CMAKE_CXX_COMPILER_LAUNCHER:${CMAKE_CXX_COMPILER_LAUNCHER}")
11	endif()	11	endif()
12		12
13	project(TORCHNPU CXX C)	13	project(TORCHNPU CXX C)
14	add_compile_options(-fmacro-prefix-map=${CMAKE_SOURCE_DIR}/=)	14	add_compile_options(-fmacro-prefix-map=${CMAKE_SOURCE_DIR}/=)
15		15
16	# GCC 11.2 (used in v2.7.1) does not support -fuse-ld=mold	16	# GCC 11.2 (used in v2.7.1) does not support -fuse-ld=mold
17	find_program(MOLD_LINKER mold)	17	find_program(MOLD_LINKER mold)
18	if(MOLD_LINKER)	18	if(MOLD_LINKER)
19	set(MOLD_WRAPPER_DIR "${CMAKE_BINARY_DIR}/mold-wrapper")	19	set(MOLD_WRAPPER_DIR "${CMAKE_BINARY_DIR}/mold-wrapper")
20	file(MAKE_DIRECTORY "${MOLD_WRAPPER_DIR}")	20	file(MAKE_DIRECTORY "${MOLD_WRAPPER_DIR}")
21	if(NOT EXISTS "${MOLD_WRAPPER_DIR}/ld")	21	if(NOT EXISTS "${MOLD_WRAPPER_DIR}/ld")
22	file(CREATE_LINK "${MOLD_LINKER}" "${MOLD_WRAPPER_DIR}/ld" SYMBOLIC)	22	file(CREATE_LINK "${MOLD_LINKER}" "${MOLD_WRAPPER_DIR}/ld" SYMBOLIC)
23	endif()	23	endif()
24	add_link_options("-B${MOLD_WRAPPER_DIR}")	24	add_link_options("-B${MOLD_WRAPPER_DIR}")
25	message(STATUS "Using mold linker via -B wrapper: ${MOLD_LINKER}")	25	message(STATUS "Using mold linker via -B wrapper: ${MOLD_LINKER}")
26	else()	26	else()
27	message(STATUS "mold linker not found, using default linker")	27	message(STATUS "mold linker not found, using default linker")
28	endif()	28	endif()
29		29
30	set(LINUX TRUE)	30	set(LINUX TRUE)
31	set(CMAKE_INSTALL_MESSAGE NEVER)	31	set(CMAKE_INSTALL_MESSAGE NEVER)
32	# set(CMAKE_VERBOSE_MAKEFILE ON)	32	# set(CMAKE_VERBOSE_MAKEFILE ON)
33	set(CMAKE_EXPORT_COMPILE_COMMANDS ON)	33	set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
34		34
35	if(DEFINED TORCH_VERSION)	35	if(DEFINED TORCH_VERSION)
36	add_definitions(-DPYTORCH_NPU_VERSION="${TORCH_VERSION}")	36	add_definitions(-DPYTORCH_NPU_VERSION="${TORCH_VERSION}")
37	endif()	37	endif()
38		38
39	set(PLUGIN_NAME torch_npu)	39	set(PLUGIN_NAME torch_npu)

	@@ -1,1101 +1,1101 @@
1	# Copyright (c) Meta Platforms, Inc. and affiliates	1	# Copyright (c) Meta Platforms, Inc. and affiliates
2	# Owner(s): ["oncall: distributed"]	2	# Owner(s): ["oncall: distributed"]
3	import os	3	import os
4		4
5	import torch	5	import torch
6	import torch.distributed._functional_collectives as funcol	6	import torch.distributed._functional_collectives as funcol
7	from torch.distributed._tensor import DTensor	7	from torch.distributed._tensor import DTensor
8	from torch.distributed.device_mesh import _mesh_resources, DeviceMesh, init_device_mesh	8	from torch.distributed.device_mesh import _mesh_resources, DeviceMesh, init_device_mesh
9	from torch.distributed.distributed_c10d import (	9	from torch.distributed.distributed_c10d import (
10	_get_default_group,	10	_get_default_group,
11	_world,	11	_world,
12	get_global_rank,	12	get_global_rank,
13	get_world_size,	13	get_world_size,
14	init_process_group,	14	init_process_group,
15	is_initialized,	15	is_initialized,
16	ProcessGroup,	16	ProcessGroup,
17	)	17	)
18	from torch.distributed.tensor._collective_utils import (	18	from torch.distributed.tensor._collective_utils import (
19	mesh_broadcast,	19	mesh_broadcast,
20	mesh_scatter,	20	mesh_scatter,
21	unpad_tensor,	21	unpad_tensor,
22	)	22	)
23	from torch.distributed.tensor.placement_types import _Partial, Shard	23	from torch.distributed.tensor.placement_types import _Partial, Shard
24	from torch.testing._internal.distributed._tensor.common_dtensor import DTensorTestBase	24	from torch.testing._internal.distributed._tensor.common_dtensor import DTensorTestBase
25	from torch.testing._internal.distributed.fake_pg import FakeStore	25	from torch.testing._internal.distributed.fake_pg import FakeStore
26		26
27	import torch_npu	27	import torch_npu
28	from torch_npu.testing.common_distributed import with_comms, skipIfUnsupportMultiNPU	28	from torch_npu.testing.common_distributed import with_comms, skipIfUnsupportMultiNPU
29	from torch_npu.testing.testcase import run_tests	29	from torch_npu.testing.testcase import run_tests
30		30
31		31
32	def _get_device_type(world_size):	32	def _get_device_type(world_size):
33	if (	33	if (
34	torch.npu.is_available()	34	torch.npu.is_available()
35	and torch.npu.device_count() >= world_size	35	and torch.npu.device_count() >= world_size
36	and torch.distributed.is_hccl_available()	36	and torch.distributed.is_hccl_available()
37	):	37	):
38	device_type = "npu"	38	device_type = "npu"
39	else:	39	else:

	@@ -1,57 +1,57 @@
1	#include <c10d/default_comm_hooks.hpp>	1	#include <c10d/default_comm_hooks.hpp>
2	#include <c10/core/ScalarType.h>	2	#include <c10/core/ScalarType.h>
3	#include <c10/util/Exception.h>	3	#include <c10/util/Exception.h>
4		4
5	#include <c10d/ProcessGroup.hpp>	5	#include <c10d/ProcessGroup.hpp>
6	#include <c10d/comm.hpp>	6	#include <c10d/comm.hpp>
7	#include <torch/torch.h>	7	#include <torch/torch.h>
8		8
9	#include "torch_npu/csrc/core/npu/NPUException.h"	9	#include "torch_npu/csrc/core/npu/NPUException.h"
10		10
11	namespace c10d {	11	namespace c10d {
12		12
13	c10::intrusive_ptr<c10::ivalue::Future> AllReduceCommHook::runHook(	13	c10::intrusive_ptr<c10::ivalue::Future> AllReduceCommHook::runHook(
14	GradBucket& bucket)	14	GradBucket& bucket)
15	{	15	{
16	std::vector<at::Tensor> tensors = {bucket.getBufferRef()};	16	std::vector<at::Tensor> tensors = {bucket.getBufferRef()};
17	// Apply the division first to avoid overflow, especially for FP16.	17	// Apply the division first to avoid overflow, especially for FP16.
18	tensors[0] /= state_->getSize();	18	tensors[0] /= state_->getSize();
19	return state_->allreduce(tensors)->getFuture();	19	return state_->allreduce(tensors)->getFuture();
20	}	20	}
21		21
22	c10::intrusive_ptr<c10::ivalue::Future> FP16CompressCommHook::runHook(	22	c10::intrusive_ptr<c10::ivalue::Future> FP16CompressCommHook::runHook(
23	GradBucket& bucket)	23	GradBucket& bucket)
24	{	24	{
25	auto compressed_tensor = bucket.getBufferRef().to(torch::kFloat16);	25	auto compressed_tensor = bucket.getBufferRef().to(torch::kFloat16);
26	// Apply the division first to avoid overflow.	26	// Apply the division first to avoid overflow.
27	compressed_tensor /= state_->getSize();	27	compressed_tensor /= state_->getSize();
28	std::vector<at::Tensor> tensors = {compressed_tensor};	28	std::vector<at::Tensor> tensors = {compressed_tensor};
29		29
30	auto allreduce_fut = state_->allreduce(tensors)->getFuture();	30	auto allreduce_fut = state_->allreduce(tensors)->getFuture();
31	auto decompressed_tensor = bucket.getBufferRef();	31	auto decompressed_tensor = bucket.getBufferRef();
32	auto decompress = [decompressed_tensor](c10::ivalue::Future& allreduce_fut) {	32	auto decompress = [decompressed_tensor](c10::ivalue::Future& allreduce_fut) {
33	auto result = allreduce_fut.value();	33	auto result = allreduce_fut.value();
34	TORCH_INTERNAL_ASSERT(	34	TORCH_INTERNAL_ASSERT(
35	result.isTensorList(),	35	result.isTensorList(),
36	"ProcessGroup::allreduce should return TensorList", DIST_ERROR(ErrCode::INTERNAL));	36	"ProcessGroup::allreduce should return TensorList", DIST_ERROR(ErrCode::INTERNAL));
37		37
38	auto reduce_tensor = result.toTensorVector()[0];	38	auto reduce_tensor = result.toTensorVector()[0];
39	TORCH_INTERNAL_ASSERT_DEBUG_ONLY(	39	TORCH_INTERNAL_ASSERT_DEBUG_ONLY(
40	reduce_tensor.scalar_type() == at::ScalarType::Half,	40	reduce_tensor.scalar_type() == at::ScalarType::Half,
41	"Expected reduced tensor to be fp16 in FP16CompressHook, but got type ",	41	"Expected reduced tensor to be fp16 in FP16CompressHook, but got type ",
42	reduce_tensor.scalar_type(), DIST_ERROR(ErrCode::TYPE)	42	reduce_tensor.scalar_type(), DIST_ERROR(ErrCode::TYPE)
43	);	43	);
44	decompressed_tensor.copy_(reduce_tensor);	44	decompressed_tensor.copy_(reduce_tensor);
45	return c10::IValue(decompressed_tensor);	45	return c10::IValue(decompressed_tensor);
46	};	46	};
47		47
48	return allreduce_fut->then(decompress, allreduce_fut->elementType());	48	return allreduce_fut->then(decompress, allreduce_fut->elementType());
49	}	49	}
50		50
51	c10::intrusive_ptr<c10::ivalue::Future> _AllReduceBySumCommHook::runHook(GradBucket& bucket)	51	c10::intrusive_ptr<c10::ivalue::Future> _AllReduceBySumCommHook::runHook(GradBucket& bucket)
52	{	52	{
53	std::vector<at::Tensor> tensors = {bucket.getBufferRef()};	53	std::vector<at::Tensor> tensors = {bucket.getBufferRef()};
54	return state_->allreduce(tensors)->getFuture();	54	return state_->allreduce(tensors)->getFuture();
55	}	55	}
56		56
57	} // namespace c10d	57	} // namespace c10d

	@@ -1,48 +1,48 @@
1	#pragma once	1	#pragma once
2		2
3	#include <c10d/ProcessGroup.hpp>	3	#include <c10d/ProcessGroup.hpp>
4	#include <c10d/comm.hpp>	4	#include <c10d/comm.hpp>
5		5
6	namespace c10d {	6	namespace c10d {
7		7
8	enum class BuiltinCommHookType {	8	enum class BuiltinCommHookType {
9	ALLREDUCE = 1,	9	ALLREDUCE = 1,
10	FP16_COMPRESS = 2,	10	FP16_COMPRESS = 2,
11	};	11	};
12		12
13	class AllReduceCommHook : public CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>> {	13	class AllReduceCommHook : public CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>> {
14	public:	14	public:
15	explicit AllReduceCommHook(c10::intrusive_ptr<ProcessGroup> state)	15	explicit AllReduceCommHook(c10::intrusive_ptr<ProcessGroup> state)
16	: CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>>(state) {}	16	: CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>>(state) {}
17		17
18	~AllReduceCommHook() override = default;	18	~AllReduceCommHook() override = default;
19		19
20	c10::intrusive_ptr<c10::ivalue::Future> runHook(GradBucket& bucket) override;	20	c10::intrusive_ptr<c10::ivalue::Future> runHook(GradBucket& bucket) override;
21	};	21	};
22		22
23	class FP16CompressCommHook : public CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>> {	23	class FP16CompressCommHook : public CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>> {
24	public:	24	public:
25	explicit FP16CompressCommHook(c10::intrusive_ptr<ProcessGroup> state)	25	explicit FP16CompressCommHook(c10::intrusive_ptr<ProcessGroup> state)
26	: CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>>(state) {}	26	: CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>>(state) {}
27		27
28	~FP16CompressCommHook() override = default;	28	~FP16CompressCommHook() override = default;
29		29
30	c10::intrusive_ptr<c10::ivalue::Future> runHook(GradBucket& bucket) override;	30	c10::intrusive_ptr<c10::ivalue::Future> runHook(GradBucket& bucket) override;
31	};	31	};
32		32
33	// Almost same as AllReduceCommHook, but without division inside the hook.	33	// Almost same as AllReduceCommHook, but without division inside the hook.
34	// This enables the optimization of fusing copy and division and saves one scan	34	// This enables the optimization of fusing copy and division and saves one scan
35	// over all the input parameters, when no communication hook is provided by the user.	35	// over all the input parameters, when no communication hook is provided by the user.
36	// Only used internally and not released as a public built-in communication hook.	36	// Only used internally and not released as a public built-in communication hook.
37	class _AllReduceBySumCommHook	37	class _AllReduceBySumCommHook
38	: public CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>> {	38	: public CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>> {
39	public:	39	public:
40	explicit _AllReduceBySumCommHook(c10::intrusive_ptr<ProcessGroup> state)	40	explicit _AllReduceBySumCommHook(c10::intrusive_ptr<ProcessGroup> state)
41	: CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>>(state) {}	41	: CppCommHookInterface<c10::intrusive_ptr<ProcessGroup>>(state) {}
42		42
43	~_AllReduceBySumCommHook() override = default;	43	~_AllReduceBySumCommHook() override = default;
44		44
45	c10::intrusive_ptr<c10::ivalue::Future> runHook(GradBucket& bucket) override;	45	c10::intrusive_ptr<c10::ivalue::Future> runHook(GradBucket& bucket) override;
46	};	46	};
47		47
48	} // namespace c10d	48	} // namespace c10d

	@@ -1,100 +1,99 @@
1	import torch	1	import torch
2	from torch._dynamo.variables import TorchInGraphFunctionVariable	2	from torch._dynamo.variables import TorchInGraphFunctionVariable
3	from torch._dynamo.trace_rules import manual_torch_name_rule_map, SkipFunctionVariable	3	from torch._dynamo.trace_rules import manual_torch_name_rule_map, SkipFunctionVariable
4	import torch._dynamo.variables.torch as torch_module	4	import torch._dynamo.variables.torch as torch_module
5	from torch._dynamo.utils import common_constant_types	5	from torch._dynamo.utils import common_constant_types
6	import torch_npu	6	import torch_npu
7		7
8	__all__ = []	8	__all__ = []
9		9
10	torch_non_c_binding_in_graph_functions_npu = dict.fromkeys(	10	torch_non_c_binding_in_graph_functions_npu = dict.fromkeys(
11	[	11	[
12	"torch.npu.current_stream",	12	"torch.npu.current_stream",
13	"torch.npu.default_stream",	13	"torch.npu.default_stream",
14	"torch.npu.stream",	14	"torch.npu.stream",
15	"torch.npu.set_stream",	15	"torch.npu.set_stream",
16	"torch_npu.npu.utils.synchronize",	16	"torch_npu.npu.utils.synchronize",
17	"torch.npu.current_device",	17	"torch.npu.current_device",
18	"torch.npu.get_device_capability",	18	"torch.npu.get_device_capability",
19	"torch.npu.get_device_properties",	19	"torch.npu.get_device_properties",
20	"torch.npu.graphs.graph_pool_handle",	20	"torch.npu.graphs.graph_pool_handle",
21	"torch.npu.ipc_collect",	21	"torch.npu.ipc_collect",
22	"torch.npu.is_available",	22	"torch.npu.is_available",
23	"torch.npu.memory._dump_snapshot",	23	"torch.npu.memory._dump_snapshot",
24	"torch.npu.memory._free_mutex",	24	"torch.npu.memory._free_mutex",
25	"torch.npu.memory._record_memory_history_impl",	25	"torch.npu.memory._record_memory_history_impl",
26	"torch.npu.memory._set_allocator_settings",	26	"torch.npu.memory._set_allocator_settings",
27	"torch.npu.memory.empty_cache",	27	"torch.npu.memory.empty_cache",
28	"torch.npu.mem_get_info",	28	"torch.npu.mem_get_info",
29	"torch.npu.memory.reset_accumulated_host_memory_stats",	29	"torch.npu.memory.reset_accumulated_host_memory_stats",
30	"torch.npu.memory.reset_accumulated_memory_stats",	30	"torch.npu.memory.reset_accumulated_memory_stats",
31	"torch.npu.memory.reset_max_memory_allocated",	31	"torch.npu.memory.reset_max_memory_allocated",
32	"torch.npu.memory.reset_max_memory_cached",	32	"torch.npu.memory.reset_max_memory_cached",
33	"torch.npu.memory.reset_peak_host_memory_stats",	33	"torch.npu.memory.reset_peak_host_memory_stats",
34	"torch.npu.memory.reset_peak_memory_stats",	34	"torch.npu.memory.reset_peak_memory_stats",
35	"torch.npu.memory.get_per_process_memory_fraction",	35	"torch.npu.memory.get_per_process_memory_fraction",
36	"torch.npu.memory.set_per_process_memory_fraction",	36	"torch.npu.memory.set_per_process_memory_fraction",
37	"torch.npu.random.manual_seed_all",	37	"torch.npu.random.manual_seed_all",
38	"torch.npu.random.manual_seed",	38	"torch.npu.random.manual_seed",
39	"torch.npu.random.seed_all",	39	"torch.npu.random.seed_all",
40	"torch.npu.random.seed",	40	"torch.npu.random.seed",
41	"torch.npu.set_sync_debug_mode",	41	"torch.npu.set_sync_debug_mode",
42	"torch.npu._set_rng_state_offset",	42	"torch.npu._set_rng_state_offset",
43	"torch.npu._get_generator",	43	"torch.npu._get_generator",
44	"torch.npu._memory_viz._frames_fmt",	44	"torch.npu._memory_viz._frames_fmt",
45	"torch.npu._memory_viz._frame_fmt",	45	"torch.npu._memory_viz._frame_fmt",
46	"torch.npu.amp.autocast_mode.custom_bwd",	46	"torch.npu.amp.autocast_mode.custom_bwd",
47	"torch.npu.amp.autocast_mode.custom_fwd",	47	"torch.npu.amp.autocast_mode.custom_fwd",
48	"torch.npu.is_initialized",	48	"torch.npu.is_initialized",
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	TorchInGraphFunctionVariable,	53	TorchInGraphFunctionVariable,
54	)	54	)
55		55
56	torch_c_binding_in_graph_functions_npu = dict.fromkeys(	56	torch_c_binding_in_graph_functions_npu = dict.fromkeys(
57	[	57	[
58	"torch_npu._C._npu_changeCurrentAllocator",	58	"torch_npu._C._npu_changeCurrentAllocator",
59	"torch_npu._C._npu_npuCachingAllocator_set_allocator_settings",	59	"torch_npu._C._npu_npuCachingAllocator_set_allocator_settings",
60	"torch_npu._C._npu_emptyCache",	60	"torch_npu._C._npu_emptyCache",
61	"torch_npu._C._npu_getAllocator",	61	"torch_npu._C._npu_getAllocator",
62	"torch_npu._C._npu_getCheckpointState",	62	"torch_npu._C._npu_getCheckpointState",
63	"torch_npu._C._npu_getCurrentStream",	63	"torch_npu._C._npu_getCurrentStream",
64	"torch_npu._C._npu_getDefaultStream",	64	"torch_npu._C._npu_getDefaultStream",
65	"torch_npu._C._npu_init",	65	"torch_npu._C._npu_init",
66	"torch_npu._C._npu_ipc_collect",	66	"torch_npu._C._npu_ipc_collect",
67	"torch_npu._C._npu_resetAccumulatedHostMemoryStats",	67	"torch_npu._C._npu_resetAccumulatedHostMemoryStats",
68	"torch_npu._C._npu_resetPeakHostMemoryStats",	68	"torch_npu._C._npu_resetPeakHostMemoryStats",
69	"torch_npu._C._npu_resetPeakMemoryStats",	69	"torch_npu._C._npu_resetPeakMemoryStats",
70	"torch_npu._C._npu_set_sync_debug_mode",	70	"torch_npu._C._npu_set_sync_debug_mode",
71	"torch_npu._C._npu_setDevice",	71	"torch_npu._C._npu_setDevice",
72	"torch_npu._C._npu_getMemoryFraction",	72	"torch_npu._C._npu_getMemoryFraction",
73	"torch_npu._C._npu_setMemoryFraction",	73	"torch_npu._C._npu_setMemoryFraction",
74	"torch_npu._C._npu_synchronize",	74	"torch_npu._C._npu_synchronize",
75	"torch_npu._C._npu_resetAccumulatedMemoryStats",	75	"torch_npu._C._npu_resetAccumulatedMemoryStats",
76	"torch_npu._C._npu_hasPrimaryContext",	76	"torch_npu._C._npu_hasPrimaryContext",
77	"torch_npu._C._npu_setStream",	77	"torch_npu._C._npu_setStream",
78	],	78	],
79	TorchInGraphFunctionVariable,	79	TorchInGraphFunctionVariable,
80	)	80	)
81		81
82	skip_functions_npu = dict.fromkeys(	82	skip_functions_npu = dict.fromkeys(
83	[	83	[
84	"torch.npu.set_device",	84	"torch.npu.set_device",
85	],	85	],
86	SkipFunctionVariable	86	SkipFunctionVariable
87	)	87	)
88		88
89		89
90	def _patch_npu_trace_rules():	90	def _patch_npu_trace_rules():
91	torch._dynamo.trace_rules.clear_lru_cache()	91	torch._dynamo.trace_rules.clear_lru_cache()
92	torch._dynamo.trace_rules.torch_name_rule_map.append(torch_non_c_binding_in_graph_functions_npu)	92	torch._dynamo.trace_rules.torch_name_rule_map.append(torch_non_c_binding_in_graph_functions_npu)
93	torch._dynamo.trace_rules.torch_name_rule_map.append(torch_c_binding_in_graph_functions_npu)	93	torch._dynamo.trace_rules.torch_name_rule_map.append(torch_c_binding_in_graph_functions_npu)
94	torch._dynamo.trace_rules.torch_name_rule_map.append(skip_functions_npu)	94	torch._dynamo.trace_rules.torch_name_rule_map.append(skip_functions_npu)
95	torch_module.constant_fold_functions[torch.npu.current_device] = True	95	torch_module.constant_fold_functions[torch.npu.current_device] = True
96	torch_module.constant_fold_functions[torch.npu.get_device_properties] = True	96	torch_module.constant_fold_functions[torch.npu.get_device_properties] = True
97	torch_module.constant_fold_functions_need_guards[torch.npu.current_device] = True	97	torch_module.constant_fold_functions_need_guards[torch.npu.current_device] = True
98	torch_module.constant_fold_functions[torch.npu.is_available] = True	98	torch_module.constant_fold_functions[torch.npu.is_available] = True
99	common_constant_types.add(torch_npu._C._NPUDeviceProperties)	99	common_constant_types.add(torch_npu._C._NPUDeviceProperties)
100