feat: 新增 --clear-l0 选项，初始化 L0A/L0B/L0C-ops-test-kit-AtomGit

M.claude/skills/ttk-how-run-test/references/kernel-params.md+1-0

	@@ -57,6 +57,7 @@ python3 -m ttk kernel -i cases.csv --co
57	\| `--clear-atomic` \| 强制在算子执行前清零输出和 workspace \| 关闭 \|	57	\| `--clear-atomic` \| 强制在算子执行前清零输出和 workspace \| 关闭 \|
58	\| `--clear-ub` \| 执行前将 UB 填充为指定值（默认清零） \| 关闭 \|	58	\| `--clear-ub` \| 执行前将 UB 填充为指定值（默认清零） \| 关闭 \|
59	\| `--clear-l1` \| 执行前将 L1 填充为指定值（默认清零） \| 关闭 \|	59	\| `--clear-l1` \| 执行前将 L1 填充为指定值（默认清零） \| 关闭 \|
		60	+\| `--clear-l0` \| 执行前将 L0A/L0B/L0C 初始化为指定值（默认清零；L0A/L0B 填充指定值，L0C 为 matmul 计算结果，值为 0 时全零） \| 关闭 \|
60	\| `--simt-ub` \| SIMT 模式 UB 大小 \| 无 \|	61	\| `--simt-ub` \| SIMT 模式 UB 大小 \| 无 \|
61	\| `--simt-stack-dcu` \| SIMT 模式 DCU 栈大小 \| 无 \|	62	\| `--simt-stack-dcu` \| SIMT 模式 DCU 栈大小 \| 无 \|
62	\| `--simt-stack-dvg` \| SIMT 模式 DVG 栈大小 \| 无 \|	63	\| `--simt-stack-dvg` \| SIMT 模式 DVG 栈大小 \| 无 \|

Mdocs/Operator_Test_Guides/Kernel_Test_Guide.md+1-0

	@@ -147,6 +147,7 @@ python3 -m ttk kernel -i add.csv --co
147	\| `--clear-atomic` \| \| 强制在算子执行前清零输出和workspace \| 关闭 \|	147	\| `--clear-atomic` \| \| 强制在算子执行前清零输出和workspace \| 关闭 \|
148	\| `--clear-ub` \| \| 执行前将UB填充为指定值（默认清零） \| 关闭 \|	148	\| `--clear-ub` \| \| 执行前将UB填充为指定值（默认清零） \| 关闭 \|
149	\| `--clear-l1` \| \| 执行前将L1填充为指定值（默认清零） \| 关闭 \|	149	\| `--clear-l1` \| \| 执行前将L1填充为指定值（默认清零） \| 关闭 \|
		150	+\| `--clear-l0` \| \| 执行前将L0A/L0B/L0C初始化为指定值（默认清零；L0A/L0B填充指定值，L0C为matmul计算结果，值为0时全零） \| 关闭 \|
150	\| `--simt-ub` \| \| SIMT 模式 UB 大小 \| 无 \|	151	\| `--simt-ub` \| \| SIMT 模式 UB 大小 \| 无 \|
151	\| `--simt-stack-dcu` \| \| SIMT 模式 DCU 栈大小 \| 无 \|	152	\| `--simt-stack-dcu` \| \| SIMT 模式 DCU 栈大小 \| 无 \|
152	\| `--force-block-dim` \| \| 强制指定 block_dim \| 无 \|	153	\| `--force-block-dim` \| \| 强制指定 block_dim \| 无 \|

Mtests/cli/test_manual_data_args.py+10-8

	@@ -3,12 +3,12 @@
3	# This program is free software, you can redistribute it and/or modify it under the terms and conditions of	3	# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	# CANN Open Software License Agreement Version 2.0 (the "License").	4	# CANN Open Software License Agreement Version 2.0 (the "License").
5	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
6	-# THIS FILE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
9	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
10	"""manual-data 两阶段模式测试：--manual-data-dirs / --no-prof 参数解析、	10	"""manual-data 两阶段模式测试：--manual-data-dirs / --no-prof 参数解析、
11	-prepare/replay 模式分发、互斥校验、--clear-ub/--clear-l1 数值解析。"""	11	+prepare/replay 模式分发、互斥校验、--clear-ub/--clear-l1/--clear-l0 数值解析。"""
12		12
13	import argparse	13	import argparse
14	import logging	14	import logging
	@@ -152,7 +152,7 @@ def test_prepare_accepts_explicit_output_directory(tmp_path):
152		152
153		153
154	@pytest.mark.parametrize(	154	@pytest.mark.parametrize(
155	- "mutate, message",	155	+ ("mutate", "message"),
156	[	156	[
157	# output/full dump 不能在设备执行前生成	157	# output/full dump 不能在设备执行前生成
158	(lambda sw: sw.dump_config.enable_output(), "--dump in,golden or --dump in"),	158	(lambda sw: sw.dump_config.enable_output(), "--dump in,golden or --dump in"),
	@@ -269,17 +269,17 @@ def test_manual_data_fields_survive_worker_pickle(tmp_path):
269	switches.manual_data_mode = "replay"	269	switches.manual_data_mode = "replay"
270	switches.manual_data_dirs = (str(tmp_path),)	270	switches.manual_data_dirs = (str(tmp_path),)
271		271
272	- restored = pickle.loads(pickle.dumps(switches))	272	+ restored = pickle.loads(pickle.dumps(switches)) # noqa: S301
273		273
274	assert restored.manual_data_mode == "replay"	274	assert restored.manual_data_mode == "replay"
275	assert restored.manual_data_dirs == (str(tmp_path),)	275	assert restored.manual_data_dirs == (str(tmp_path),)
276		276
277		277
278	-# -- --clear-ub / --clear-l1 数值解析 ----------------------------------------	278	+# -- --clear-ub / --clear-l1 / --clear-l0 数值解析 ----------------------------------------
279		279
280		280
281	@pytest.mark.parametrize(	281	@pytest.mark.parametrize(
282	- "value, expected_type, expected",	282	+ ("value", "expected_type", "expected"),
283	[	283	[
284	("7", np.int32, 7),	284	("7", np.int32, 7),
285	("0xff", np.int32, 255),	285	("0xff", np.int32, 255),
	@@ -290,7 +290,7 @@ def test_manual_data_fields_survive_worker_pickle(tmp_path):
290	ids=["int", "hex", "float", "typed-float16", "typed-uint8"],	290	ids=["int", "hex", "float", "typed-float16", "typed-uint8"],
291	)	291	)
292	def test_clear_value_parser_accepts_numeric_literals(value, expected_type, expected):	292	def test_clear_value_parser_accepts_numeric_literals(value, expected_type, expected):
293	- """--clear-ub/--clear-l1 解析十进制/十六进制/浮点/带 dtype 前缀的数值字面量。"""	293	+ """--clear-ub/--clear-l1/--clear-l0 解析十进制/十六进制/浮点/带 dtype 前缀的数值字面量。"""
294	parsed = _parse_clean_val("UB", value)	294	parsed = _parse_clean_val("UB", value)
295		295
296	assert isinstance(parsed, expected_type)	296	assert isinstance(parsed, expected_type)
	@@ -298,9 +298,11 @@ def test_clear_value_parser_accepts_numeric_literals(value, expected_type, expec
298		298
299		299
300	def test_clear_value_parser_accepts_inf_and_nan():	300	def test_clear_value_parser_accepts_inf_and_nan():
301	- """--clear-ub/--clear-l1 解析特殊浮点值 inf/nan。"""	301	+ """--clear-ub/--clear-l1/--clear-l0 解析特殊浮点值 inf/nan。"""
302	assert np.isinf(_parse_clean_val("L1", "float32(inf)"))	302	assert np.isinf(_parse_clean_val("L1", "float32(inf)"))
303	assert np.isnan(_parse_clean_val("L1", "nan"))	303	assert np.isnan(_parse_clean_val("L1", "nan"))
		304	+ assert np.isinf(_parse_clean_val("L0", "float32(inf)"))
		305	+ assert np.isnan(_parse_clean_val("L0", "nan"))
304		306
305		307
306	def test_clear_value_parser_rejects_code_injection():	308	def test_clear_value_parser_rejects_code_injection():

Mttk/cli/bridge.py+11-3

	@@ -1,3 +1,10 @@
		1	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
		2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
		3	+# CANN Open Software License Agreement Version 2.0 (the "License").
		4	+# Please refer to the License for details. You may not use this file except in compliance with the License.
		5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
		6	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
		7	+# See LICENSE in the root of the software repository for the full text of the License.
1	import logging	8	import logging
2	import pathlib	9	import pathlib
3	import re	10	import re
	@@ -187,6 +194,8 @@ def apply_kernel_args(sw, args):
187	sw.force_clear_ub = _parse_clean_val("UB", args.clear_ub)	194	sw.force_clear_ub = _parse_clean_val("UB", args.clear_ub)
188	if hasattr(args, "clear_l1") and args.clear_l1 is not None:	195	if hasattr(args, "clear_l1") and args.clear_l1 is not None:
189	sw.force_clear_l1 = _parse_clean_val("L1", args.clear_l1)	196	sw.force_clear_l1 = _parse_clean_val("L1", args.clear_l1)
		197	+ if hasattr(args, "clear_l0") and args.clear_l0 is not None:
		198	+ sw.force_clear_l0 = _parse_clean_val("L0", args.clear_l0)
190	if hasattr(args, "force_block_dim") and args.force_block_dim is not None:	199	if hasattr(args, "force_block_dim") and args.force_block_dim is not None:
191	bd = args.force_block_dim	200	bd = args.force_block_dim
192	if isinstance(bd, int):	201	if isinstance(bd, int):
	@@ -457,9 +466,8 @@ def _apply_cce(sw, value):
457	sw.dyn_switches.realtime = False	466	sw.dyn_switches.realtime = False
458	elif mode in ("c", "cst", "const"):	467	elif mode in ("c", "cst", "const"):
459	sw.cst_switches.realtime = False	468	sw.cst_switches.realtime = False
460	- elif mode in ("b", "bin", "binary"):	469	+ elif mode in ("b", "bin", "binary") and sw.bin_switches.realtime != "release":
461	- if sw.bin_switches.realtime != "release":	470	+ sw.bin_switches.realtime = False
462	- sw.bin_switches.realtime = False
463		471
464		472
465	def _apply_clear_atomic(sw, value):	473	def _apply_clear_atomic(sw, value):

Mttk/cli/kernel.py+14-0

	@@ -1,3 +1,10 @@
		1	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
		2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
		3	+# CANN Open Software License Agreement Version 2.0 (the "License").
		4	+# Please refer to the License for details. You may not use this file except in compliance with the License.
		5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
		6	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
		7	+# See LICENSE in the root of the software repository for the full text of the License.
1	from ttk.cli.bridge import (	8	from ttk.cli.bridge import (
2	apply_kernel_args,	9	apply_kernel_args,
3	args_to_switches,	10	args_to_switches,
	@@ -77,6 +84,13 @@ def _add_kernel_args(parser):
77	parser.add_argument(	84	parser.add_argument(
78	"--clear-l1", dest="clear_l1", default=None, help="Clear L1 to specified value before execution (default: 0)"	85	"--clear-l1", dest="clear_l1", default=None, help="Clear L1 to specified value before execution (default: 0)"
79	)	86	)
		87	+ parser.add_argument(
		88	+ "--clear-l0",
		89	+ dest="clear_l0",
		90	+ default=None,
		91	+ help="Clear L0A/L0B/L0C to specified value before execution (default: 0). "
		92	+ "L0A/L0B are filled with the value; L0C is the matmul result (0 when value is 0).",
		93	+ )
80		94
81	parser.add_argument("--simt-ub", dest="simt_ub", default=None, help="Force SIMT UB size in bytes")	95	parser.add_argument("--simt-ub", dest="simt_ub", default=None, help="Force SIMT UB size in bytes")
82	parser.add_argument(	96	parser.add_argument(

Mttk/core_modules/aclnn/acl_interface.py+12-15

	@@ -186,6 +186,9 @@ class AclInterface:
186	def clear_l1(self, switches: "ttk.utilities.SWITCHES"):	186	def clear_l1(self, switches: "ttk.utilities.SWITCHES"):
187	self._rts_interface.clear_l1(switches)	187	self._rts_interface.clear_l1(switches)
188		188
		189	+ def clear_l0(self, switches: "ttk.utilities.SWITCHES"):
		190	+ self._rts_interface.clear_l0(switches)
		191	+
189	def clear_ub(self, switches: "ttk.utilities.SWITCHES"):	192	def clear_ub(self, switches: "ttk.utilities.SWITCHES"):
190	self._rts_interface.clear_ub(switches)	193	self._rts_interface.clear_ub(switches)
191		194
	@@ -253,10 +256,9 @@ class AclInterface:
253	return self._create_acl_tensor_from_numpy(	256	return self._create_acl_tensor_from_numpy(
254	tensor, view_format, storage_shape, view_shape_override, acl_dtype_override	257	tensor, view_format, storage_shape, view_shape_override, acl_dtype_override
255	)	258	)
256	- else:	259	+ return self._create_acl_tensor_from_torch(
257	- return self._create_acl_tensor_from_torch(	260	+ tensor, view_format, storage_shape, view_shape_override, acl_dtype_override
258	- tensor, view_format, storage_shape, view_shape_override, acl_dtype_override	261	+ )
259	- )
260		262
261	def _create_acl_tensor_from_torch(	263	def _create_acl_tensor_from_torch(
262	self,	264	self,
	@@ -392,9 +394,8 @@ class AclInterface:
392	for t in acl_tensor_ptr_lst:	394	for t in acl_tensor_ptr_lst:
393	if t is None:	395	if t is None:
394	continue	396	continue
395	- if isinstance(t, ctypes.c_void_p):	397	+ ptr = t.value if isinstance(t, ctypes.c_void_p) else t
396	- t = t.value	398	+ self._acl_tensors.remove(ptr)
397	- self._acl_tensors.remove(t)
398	return c_ptr	399	return c_ptr
399		400
400	def create_acl_scalar(self, val: Union[numpy.ndarray, "torch.Tensor"]) -> ctypes.c_void_p:	401	def create_acl_scalar(self, val: Union[numpy.ndarray, "torch.Tensor"]) -> ctypes.c_void_p:
	@@ -439,9 +440,8 @@ class AclInterface:
439	for s in acl_scalar_ptr_lst:	440	for s in acl_scalar_ptr_lst:
440	if s is None:	441	if s is None:
441	continue	442	continue
442	- if isinstance(s, ctypes.c_void_p):	443	+ ptr = s.value if isinstance(s, ctypes.c_void_p) else s
443	- s = s.value	444	+ self._acl_scalars.remove(ptr)
444	- self._acl_scalars.remove(s)
445	return c_ptr	445	return c_ptr
446		446
447	def create_acl_array(self, val_lst: Union[tuple, list], typ: str) -> ctypes.c_void_p:	447	def create_acl_array(self, val_lst: Union[tuple, list], typ: str) -> ctypes.c_void_p:
	@@ -451,10 +451,7 @@ class AclInterface:
451	cnt = len(val_lst)	451	cnt = len(val_lst)
452	c_type = self.AclArrayDict[typ]	452	c_type = self.AclArrayDict[typ]
453	c_size = ctypes.c_uint64(cnt)	453	c_size = ctypes.c_uint64(cnt)
454	- if cnt == 0:	454	+ c_value = None if cnt == 0 else (c_type * cnt)(*val_lst)
455	- c_value = None
456	- else:
457	- c_value = (c_type * cnt)(*val_lst)
458	c_ptr = self._opbase_api_call_with_ptr_return(	455	c_ptr = self._opbase_api_call_with_ptr_return(
459	f"aclCreate{typ}Array", f"Args: value={val_lst}, size={cnt}", c_value, c_size	456	f"aclCreate{typ}Array", f"Args: value={val_lst}, size={cnt}", c_value, c_size
460	)	457	)
	@@ -569,7 +566,7 @@ class AclInterface:
569	self._free_acl_scalar_list(self._acl_scalar_lists.pop())	566	self._free_acl_scalar_list(self._acl_scalar_lists.pop())
570	while self._acl_scalars:	567	while self._acl_scalars:
571	self._free_acl_scalar(self._acl_scalars.pop())	568	self._free_acl_scalar(self._acl_scalars.pop())
572	- for typ in self.AclArrayDict.keys():	569	+ for typ in self.AclArrayDict:
573	array_sets = getattr(self, f"_acl_{typ.lower()}_arrays")	570	array_sets = getattr(self, f"_acl_{typ.lower()}_arrays")
574	while array_sets:	571	while array_sets:
575	self._free_acl_array(array_sets.pop(), typ)	572	self._free_acl_array(array_sets.pop(), typ)

Mttk/core_modules/npu/instance_refactor.py+15-18

	@@ -47,9 +47,7 @@ class NpuInstance(InstanceBase):
47	still needs one logical worker. That mode deliberately sets the stored	47	still needs one logical worker. That mode deliberately sets the stored
48	count to one and avoids querying DSMI hardware.	48	count to one and avoids querying DSMI hardware.
49	"""	49	"""
50	- if getattr(self.switches, "manual_data_mode", None) == "prepare":	50	+ if getattr(self.switches, "manual_data_mode", None) == "prepare" or self.switches.mode.is_model():
51	- self.switches.device_count = 1
52	- elif self.switches.mode.is_model():
53	self.switches.device_count = 1	51	self.switches.device_count = 1
54	elif self.switches.device_count <= 0:	52	elif self.switches.device_count <= 0:
55	if self.switches.compile_only:	53	if self.switches.compile_only:
	@@ -72,27 +70,25 @@ class NpuInstance(InstanceBase):
72	if self.switches.dev_plat == "AUTO":	70	if self.switches.dev_plat == "AUTO":
73	if self.switches.mode.is_model():	71	if self.switches.mode.is_model():
74	raise RuntimeError(f"Please specify your platform type with --plat in {self.switches.mode.name} mode")	72	raise RuntimeError(f"Please specify your platform type with --plat in {self.switches.mode.name} mode")
75	- else:	73	+ try:
76	- try:	74	+ self.switches.dev_plat = DSMIInterface().get_chip_info(0).get_complete_platform()
77	- self.switches.dev_plat = DSMIInterface().get_chip_info(0).get_complete_platform()	75	+ except Exception as e:
78	- except Exception as e:	76	+ if (
79	- if (	77	+ self.switches.compile_only
80	- self.switches.compile_only	78	+ or self.switches.validate_only
81	- or self.switches.validate_only	79	+ or getattr(self.switches, "manual_data_mode", None) == "prepare"
82	- or getattr(self.switches, "manual_data_mode", None) == "prepare"	80	+ ):
83	- ):	81	+ raise RuntimeError(
84	- raise RuntimeError(	82	+ "Try to get Ascend platform failed. Please specify it with option like: --plat=Ascend910A"
85	- "Try to get Ascend platform failed. Please specify it with option like: --plat=Ascend910A"	83	+ ) from e
86	- ) from e	84	+ raise
87	- else:
88	- raise
89	hw_info = get_npu_hw_info(self.switches.dev_plat)	85	hw_info = get_npu_hw_info(self.switches.dev_plat)
90	self.switches.short_soc_version = hw_info.get("short_soc_version")	86	self.switches.short_soc_version = hw_info.get("short_soc_version")
91	os.environ["TTK_FULL_SOC_VERSION"] = self.switches.dev_plat	87	os.environ["TTK_FULL_SOC_VERSION"] = self.switches.dev_plat
92	os.environ["TTK_SHORT_SOC_VERSION"] = self.switches.short_soc_version	88	os.environ["TTK_SHORT_SOC_VERSION"] = self.switches.short_soc_version
93		89
94	def setup_profile_object(self):	90	def setup_profile_object(self):
95	- params = tuple([self.task_keeper, self.mp_context])	91	+ params = (self.task_keeper, self.mp_context)
96	if "api_name" in self.case_original_headers:	92	if "api_name" in self.case_original_headers:
97	from .op_api import ApiProfileObject	93	from .op_api import ApiProfileObject
98		94
	@@ -185,4 +181,5 @@ class NpuInstance(InstanceBase):
185	def _compile_help_kernels():	181	def _compile_help_kernels():
186	Opc().compile_ub_clear()	182	Opc().compile_ub_clear()
187	Opc().compile_l1_clear()	183	Opc().compile_l1_clear()
		184	+ Opc().compile_l0_clear()
188	Opc().compile_warmup_kernel()	185	Opc().compile_warmup_kernel()

Mttk/core_modules/npu/op/rts_sequence.py+23-31

	@@ -38,22 +38,20 @@ def _process_total_cycles(results: list):
38	# kick-off UNKNOWN	38	# kick-off UNKNOWN
39	kicked_prof_task = [ele for ele in results if ele != "UNKNOWN"]	39	kicked_prof_task = [ele for ele in results if ele != "UNKNOWN"]
40	# Check if profiling result is valid	40	# Check if profiling result is valid
41	- profiling_data_valid = kicked_prof_task and all([isinstance(ele, (int, float)) for ele in kicked_prof_task])	41	+ profiling_data_valid = kicked_prof_task and all(isinstance(ele, (int, float)) for ele in kicked_prof_task)
42	logging.debug(f"HWTS Task Profiling result: {results}")	42	logging.debug(f"HWTS Task Profiling result: {results}")
43	if profiling_data_valid:	43	if profiling_data_valid:
44	kicked_prof_task.sort()	44	kicked_prof_task.sort()
45	return numpy.median(kicked_prof_task)	45	return numpy.median(kicked_prof_task)
46	- else:	46	+ task_data = tuple(map(str, results))
47	- task_data = tuple(map(str, results))	47	+ return ",".join(task_data)
48	- return ",".join(task_data)
49		48
50		49
51	def rts_profiling(device: RTSInterfaceBase, profiling_param: RTSProfilingParam):	50	def rts_profiling(device: RTSInterfaceBase, profiling_param: RTSProfilingParam):
52	if isinstance(device, RTSInterface):	51	if isinstance(device, RTSInterface):
53	online_obj = OnlineRtsProfiling(device, profiling_param)	52	online_obj = OnlineRtsProfiling(device, profiling_param)
54	return online_obj.do()	53	return online_obj.do()
55	- else:	54	+ raise RuntimeError(f"Unrecognized device type: {type(device)}")
56	- raise RuntimeError(f"Unrecognized device type: {type(device)}")
57		55
58		56
59	class OnlineRtsProfiling:	57	class OnlineRtsProfiling:
	@@ -167,14 +165,13 @@ class OnlineRtsProfiling:
167	f"RTS Register Binary failed, kernel object {kernel_full_path} does not exist or is invalid."	165	f"RTS Register Binary failed, kernel object {kernel_full_path} does not exist or is invalid."
168	)	166	)
169	return RTSProfilingResult.fail("RTS_BINARY_FAILURE")	167	return RTSProfilingResult.fail("RTS_BINARY_FAILURE")
170	- else:	168	+ logging.error(
171	- logging.error(	169	+ f"RTS Register Function failed, "
172	- f"RTS Register Function failed, "	170	+ f"Expect symbol {self._param.kernel_main_func_name} does not exist "
173	- f"Expect symbol {self._param.kernel_main_func_name} does not exist "	171	+ f"in {kernel_full_path}. "
174	- f"in {kernel_full_path}. "	172	+ f"this is usually caused by wrong tiling key"
175	- f"this is usually caused by wrong tiling key"	173	+ )
176	- )	174	+ return RTSProfilingResult.fail("RTS_FUNCTION_FAILURE")
177	- return RTSProfilingResult.fail("RTS_FUNCTION_FAILURE")
178		175
179	def _rts_kernel_sequence_v2(self, stream: Optional[ctypes.c_void_p] = None):	176	def _rts_kernel_sequence_v2(self, stream: Optional[ctypes.c_void_p] = None):
180	# Profiling Preparation	177	# Profiling Preparation
	@@ -198,8 +195,9 @@ class OnlineRtsProfiling:
198	) as profiler:	195	) as profiler:
199	for repeat_idx in range(self._run_time):	196	for repeat_idx in range(self._run_time):
200	profiler.step()	197	profiler.step()
201	- self._device.clear_l1(self._switches)	198	+ self._device.clear_l0(self._switches)
202	self._device.clear_ub(self._switches)	199	self._device.clear_ub(self._switches)
		200	+ self._device.clear_l1(self._switches)
203	# self._device.test_clear_ub(self._switches)	201	# self._device.test_clear_ub(self._switches)
204	# Prepare Memory on HBM	202	# Prepare Memory on HBM
205	self._alloc_device_memory(repeat_idx)	203	self._alloc_device_memory(repeat_idx)
	@@ -257,9 +255,8 @@ class OnlineRtsProfiling:
257	ipt_ids = [id(i) for i in self._param.flatten_input_arrays if isinstance(i, numpy.ndarray)]	255	ipt_ids = [id(i) for i in self._param.flatten_input_arrays if isinstance(i, numpy.ndarray)]
258	inplace_ids = []	256	inplace_ids = []
259	for o in self._param.flatten_output_arrays:	257	for o in self._param.flatten_output_arrays:
260	- if isinstance(o, numpy.ndarray):	258	+ if isinstance(o, numpy.ndarray) and id(o) in ipt_ids:
261	- if id(o) in ipt_ids:	259	+ inplace_ids.append(id(o))
262	- inplace_ids.append(id(o))
263	return inplace_ids	260	return inplace_ids
264		261
265	def _alloc_hbm(self):	262	def _alloc_hbm(self):
	@@ -298,9 +295,8 @@ class OnlineRtsProfiling:
298	# tensor list	295	# tensor list
299	refs = []	296	refs = []
300	for t in oa:	297	for t in oa:
301	- if isinstance(t, numpy.ndarray):	298	+ if isinstance(t, numpy.ndarray) and id(t) in inplace_ids:
302	- if id(t) in inplace_ids:	299	+ refs.append(id(t))
303	- refs.append(id(t))
304	if refs:	300	if refs:
305	# inplace dynamic tensor-list	301	# inplace dynamic tensor-list
306	if len(refs) != len(oa):	302	if len(refs) != len(oa):
	@@ -371,9 +367,8 @@ class OnlineRtsProfiling:
371	for a in arrays:	367	for a in arrays:
372	if a is None:	368	if a is None:
373	continue	369	continue
374	- else:	370	+ mem = self._nd_array_maps[id(a)]
375	- mem = self._nd_array_maps[id(a)]	371	+ self._device.copy_nparray_to_hbm_ptr(a, mem)
376	- self._device.copy_nparray_to_hbm_ptr(a, mem)
377		372
378	def _free_device_memory(self):	373	def _free_device_memory(self):
379	for _, dev_address in self._nd_array_maps.items():	374	for _, dev_address in self._nd_array_maps.items():
	@@ -615,7 +610,7 @@ class OnlineRtsProfiling:
615	pass	610	pass
616		611
617	if op_prof:	612	if op_prof:
618	- op_prof = list(x["duration"] for x in op_prof)	613	+ op_prof = [x["duration"] for x in op_prof]
619	results = [op_prof[0]]	614	results = [op_prof[0]]
620	if len(op_prof) > 1:	615	if len(op_prof) > 1:
621	results = op_prof[1:]	616	results = op_prof[1:]
	@@ -629,8 +624,8 @@ class OnlineRtsProfiling:
629	cmds = ["nm", kernel_full_path]	624	cmds = ["nm", kernel_full_path]
630	try:	625	try:
631	out = subprocess.check_output(cmds).decode("utf-8")	626	out = subprocess.check_output(cmds).decode("utf-8")
632	- except Exception:	627	+ except Exception as exc:
633	- pass	628	+ logging.debug(f"nm command failed for {kernel_full_path}: {exc}")
634	else:	629	else:
635	if out:	630	if out:
636	splits = re.split(r"[ \n]", out)	631	splits = re.split(r"[ \n]", out)
	@@ -670,8 +665,5 @@ class OnlineRtsProfiling:
670	self._device.set_task_fail_callback()	665	self._device.set_task_fail_callback()
671	self._rts_task_fail_cb_set = True	666	self._rts_task_fail_cb_set = True
672	# rtSetExceptionExtInfo	667	# rtSetExceptionExtInfo
673	- try:	668	+ with contextlib.suppress(RuntimeError):
674	self._device.set_exception_extend_info(ctypes.c_void_p(ctypes.addressof(self._task_fail_cb_invoked)))	669	self._device.set_exception_extend_info(ctypes.c_void_p(ctypes.addressof(self._task_fail_cb_invoked)))
675	- except RuntimeError:
676	- # if fail, just pass.
677	- pass

Mttk/core_modules/npu/op_api/profiling.py+2-1

	@@ -432,8 +432,9 @@ class AclOpExecutor:
432	for repeat_idx in range(self._run_time):	432	for repeat_idx in range(self._run_time):
433	if not skip_profiler and profiler:	433	if not skip_profiler and profiler:
434	profiler.step()	434	profiler.step()
435	- self._dvc.clear_l1(self._switches)	435	+ self._dvc.clear_l0(self._switches)
436	self._dvc.clear_ub(self._switches)	436	self._dvc.clear_ub(self._switches)
		437	+ self._dvc.clear_l1(self._switches)
437	logging.debug(	438	logging.debug(
438	f"[AclOpExecutor dev={self._dvc.device_id}] building phase1 params, "	439	f"[AclOpExecutor dev={self._dvc.device_id}] building phase1 params, "
439	f"group={self._ctx.attributes.get('group', 'N/A')}"	440	f"group={self._ctx.attributes.get('group', 'N/A')}"

Attk/core_modules/operator/ascendc_kernels/clear_l0_ascendc.cpp+142-0

@@ -0,0 +1,142 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	R R RuiWang_15 天前文件不要带950的命名
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*
	12	+ * clear_l0 helper kernel for David-generation Ascend chips (dav-3510, ...).
	13	+ *
	14	+ * TIK cube intrinsics (mmad/load2dv2/fixpipe) are not supported on David-
	15	+ * generation chips, so L0A/L0B/L0C are cleared via tensor_api (the same
	16	+ * primitives Blaze uses):
	17	+ * GM -> L1 -> L0A/L0B -> mad -> L0C -> fixpipe -> GM.
	18	+ *
	19	+ * Tile dimensions M_DIM / K_DIM / N_DIM are passed at compile time via -D
	20	+ * macros (computed by the host from the target chip's L0A/L0B/L0C sizes).
	21	+ * One pass covers the full L0A and L0B, and as much of L0C as a single
	22	+ * mmad allows (MK2B = L0A, KN2B = L0B, MN4B <= L0C).
	23	+ *
	24	+ * Defaults below match Ascend950 (L0A/L0B = 64KB, L0C = 256KB):
	25	+ * M=256, K=128, N=256
	26	+ * L0A tile: 256 * 128 * 2B = 64KB (full L0A)
	27	+ * L0B tile: 128 * 256 * 2B = 64KB (full L0B)
	28	+ * L0C tile: 256 * 256 * 4B = 256KB (full L0C)
	29	+ *
	30	+ * Kernel args (both from TTK helper launch):
	31	+ * input : (MK + KN) * 2B fp16, every element = clean_val
	32	+ * first MK elements act as A (M x K ND), next KN as B (K x N ND)
	33	+ * output : MN 4B fp32 dummy sink for the fixpipe result (A x B)
	34	+ * host reads it back after sync and verifies every element
	35	+ * == K * clean_val^2, proving the full
	36	+ * GM->L1->L0A/L0B->Mmad->L0C->fixpipe chain ran
	37	+ *
	38	+ * blockDim = AIC core num; every core runs the identical sequence so the
	39	+ * L0A/L0B/L0C of every core are overwritten with deterministic data.
	40	+ */
	41	+
	42	+#include "kernel_operator.h"
	43	+#include "tensor_api/tensor.h"
	44	+
	45	+namespace {
	46	+using namespace AscendC::Te;
	47	+
	48	+#ifndef M_DIM
	49	+#define M_DIM 256
	50	+#endif
	51	+#ifndef K_DIM
	52	+#define K_DIM 128
	53	+#endif
	54	+#ifndef N_DIM
	55	+#define N_DIM 256
	56	+#endif
	57	+
	58	+constexpr int32_t L1_EVENT_ID = 0;
	59	+constexpr int32_t L0_EVENT_ID = 1;
	60	+constexpr int32_t L0C_EVENT_ID = 2;
	61	+constexpr uint8_t FINAL_ACCUMULATION = 3;
	62	+} // namespace
	63	+
	64	+extern "C" __global__ __aicore__ __cube__ void clear_l0(GM_ADDR input, GM_ADDR output) {
	65	+ if ASCEND_IS_AIV {
	66	+ return;
	67	+ }
	68	+ using namespace AscendC::Te;
	69	+
	70	+ auto gmAPtr = reinterpret_cast<__gm__ half *>(input);
	71	+ auto gmBPtr = reinterpret_cast<__gm__ half >(input) + M_DIM K_DIM;
	72	+ auto gmCPtr = reinterpret_cast<__gm__ float *>(output);
	73	+
	74	+ auto gmATensor = MakeTensor(MakeMemPtr(gmAPtr), MakeFrameLayout<NDExtLayoutPtn>(M_DIM, K_DIM));
	75	+ auto gmBTensor = MakeTensor(MakeMemPtr(gmBPtr), MakeFrameLayout<NDExtLayoutPtn>(K_DIM, N_DIM));
	76	+ auto gmCTensor = MakeTensor(MakeMemPtr(gmCPtr), MakeFrameLayout<NDExtLayoutPtn>(M_DIM, N_DIM));
	77	+
	78	+ __cbuf__ half l1ABuf[M_DIM * K_DIM];
	79	+ __cbuf__ half l1BBuf[K_DIM * N_DIM];
	80	+ __ca__ half l0ABuf[M_DIM * K_DIM];
	81	+ __cb__ half l0BBuf[K_DIM * N_DIM];
	82	+ __cc__ float l0CBuf[M_DIM * N_DIM];
	83	+
	84	+ auto l1ATensor = MakeTensor(MakeMemPtr(l1ABuf), MakeFrameLayout<NZLayoutPtn, half>(M_DIM, K_DIM));
	85	+ auto l1BTensor = MakeTensor(MakeMemPtr(l1BBuf), MakeFrameLayout<NZLayoutPtn, half>(K_DIM, N_DIM));
	86	+ auto l0ATensor = MakeTensor(MakeMemPtr(l0ABuf), MakeFrameLayout<NZLayoutPtn, half>(M_DIM, K_DIM));
	87	+ auto l0BTensor = MakeTensor(MakeMemPtr(l0BBuf), MakeFrameLayout<ZNLayoutPtn, half>(K_DIM, N_DIM));
	88	+ auto l0CTensor = MakeTensor(MakeMemPtr(l0CBuf), MakeFrameLayout<NZLayoutPtn>(M_DIM, N_DIM));
	89	+
	90	+ auto copyGM2L1Atom = MakeCopy(CopyGM2L1{}, CopyGM2L1TraitDefault{});
	91	+ auto copyL12L0AAtom = MakeCopy(CopyL12L0A{}, CopyL12L0ATraitDefault{});
	92	+ auto copyL12L0BAtom = MakeCopy(CopyL12L0B{}, CopyL12L0BTraitDefault{});
	93	+ auto copyL0C2GMAtom = MakeCopy(CopyL0C2GM{}, CopyL0C2GMTraitDefault{});
	94	+ auto mmadAtom = MakeMmad(MmadOperation{}, MmadTraitDefault{});
	95	+
	96	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(L1_EVENT_ID);
	97	+ AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(L0_EVENT_ID);
	98	+ AscendC::SetFlag<AscendC::HardEvent::FIX_M>(L0C_EVENT_ID);
	99	+
	100	+ // GM -> L1
	101	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(L1_EVENT_ID);
	102	+ Copy(copyGM2L1Atom, l1ATensor, gmATensor);
	103	+ Copy(copyGM2L1Atom, l1BTensor, gmBTensor);
	104	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(L1_EVENT_ID);
	105	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(L1_EVENT_ID);
	106	+
	107	+ // L1 -> L0A/L0B
	108	+ AscendC::WaitFlag<AscendC::HardEvent::M_MTE1>(L0_EVENT_ID);
	109	+ Copy(copyL12L0AAtom, l0ATensor, l1ATensor);
	110	+ Copy(copyL12L0BAtom, l0BTensor, l1BTensor);
	111	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(L1_EVENT_ID);
	112	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(L0_EVENT_ID);
	113	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(L0_EVENT_ID);
	114	+
	115	+ // L0A x L0B -> L0C
	116	+ AscendC::WaitFlag<AscendC::HardEvent::FIX_M>(L0C_EVENT_ID);
	117	+ MmadParams mmadParams;
	118	+ mmadParams.m = M_DIM;
	119	+ mmadParams.n = N_DIM;
	120	+ mmadParams.k = K_DIM;
	121	+ mmadParams.cmatrixInitVal = true;
	122	+ mmadParams.unitFlag = FINAL_ACCUMULATION;
	123	+ Mmad(mmadAtom.with(mmadParams), l0CTensor, l0ATensor, l0BTensor);
	124	+ AscendC::SetFlag<AscendC::HardEvent::M_FIX>(L0C_EVENT_ID);
	125	+ AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(L0_EVENT_ID);
	126	+
	127	+ // L0C -> GM
	128	+ AscendC::WaitFlag<AscendC::HardEvent::M_FIX>(L0C_EVENT_ID);
	129	+ FixpipeParams fixpipeParams;
	130	+ fixpipeParams.unitFlag = FINAL_ACCUMULATION;
	131	+ Copy(copyL0C2GMAtom.with(fixpipeParams), gmCTensor, l0CTensor);
	132	+ AscendC::SetFlag<AscendC::HardEvent::FIX_M>(L0C_EVENT_ID);
	133	+
	134	+ AscendC::WaitFlag<AscendC::HardEvent::M_MTE1>(L0_EVENT_ID);
	135	+ AscendC::WaitFlag<AscendC::HardEvent::FIX_M>(L0C_EVENT_ID);
	136	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(L1_EVENT_ID);
	137	+
	138	+ // Info-level log: only printed in CPU debug mode; compiled out (empty macro)
	139	+ // on real-device builds, so it never floods stdout on board.
	140	+ KERNEL_LOG(KERNEL_INFO, "[clear_l0] core=%u l0 injected, cycle=%lu",
	141	+ static_cast<uint32_t>(AscendC::GetBlockIdx()), AscendC::GetSystemCycle());
	142	+}

Mttk/core_modules/operator/helper_kernels.py+78-3

	@@ -38,9 +38,7 @@ def clear_ub(ipt: dict, full_soc_version: str, core_type: str, kernel_name: str
38	return tik_instance	38	return tik_instance
39		39
40		40
41	-def test_clear_ub(	41	+def test_clear_ub(output, full_soc_version: str, core_type: str, clean_val: numpy.generic, kernel_name: str):
42	- output, full_soc_version: str, core_type: str, clean_val: numpy.generic, kernel_name: str = "test_clear_ub"
43	-):
44	"""Test UB data after clear"""	42	"""Test UB data after clear"""
45	tbe_platform.set_current_compile_soc_info(full_soc_version, core_type)	43	tbe_platform.set_current_compile_soc_info(full_soc_version, core_type)
46	tik_instance = tik.Tik()	44	tik_instance = tik.Tik()
	@@ -123,6 +121,83 @@ def clear_l1(ipt: dict, full_soc_version: str, kernel_name: str = "clear_l1"):
123	return tik_instance	121	return tik_instance
124		122
125		123
		124	+def clear_l0(ipt: dict, full_soc_version: str, kernel_name: str = "clear_l0"):
		125	+ """
		126	+ clear L0A/L0B/L0C data via two-step matmul (GM->L1->L0A/L0B->L0C).
		127	+ matmul internally uses mmad instruction to load L1->L0A/L0B and write L0C.
		128	+ Step1 fills L0A fully and writes L0C; Step2 fills L0B fully and overwrites L0C.
		129	+ On Ascend950, matmul uses f16f16f16 (L0C as float16) with api_check_support
		130	+ patched because tik_api_map lacks the 950 entry.
		131	+ On Ascend910B, matmul uses f16f16f32 (L0C as float32).
		132	+ """
		133	+ import tbe.common.platform.platform_info as _pi
		134	+
		135	+ _orig_check = _pi.api_check_support
		136	+
		137	+ def _patched_check(name, dtype_str):
		138	+ if name == "tik.matmul" and dtype_str == "f16f16f16":
		139	+ return True
		140	+ if name == "tik.load2dv2" and dtype_str == "float16":
		141	+ return True
		142	+ return _orig_check(name, dtype_str)
		143	+
		144	+ _pi.api_check_support = _patched_check
		145	+ import tbe.tik.api.cube.matmul as _matmul_mod
		146	+
		147	+ _matmul_mod.api_check_support = _patched_check
		148	+ import tbe.tik.tik_lib.tik_mmad_convert_api.tik_mmad_convert_operation as _load2d_mod
		149	+
		150	+ _load2d_mod.api_check_support = _patched_check
		151	+
		152	+ tbe_platform.set_current_compile_soc_info(full_soc_version, "AiCore")
		153	+ tik_instance = tik.Tik()
		154	+ core_num = tbe_platform.get_soc_spec("CORE_NUM")
		155	+
		156	+ l0a_size = tbe_platform.get_soc_spec(tbe_platform.L0A_SIZE)
		157	+ l0b_size = tbe_platform.get_soc_spec(tbe_platform.L0B_SIZE)
		158	+ dtype = "float16"
		159	+ dtype_bytes = numpy.dtype(dtype).itemsize
		160	+
		161	+ short_soc = full_soc_version
		162	+ is_950 = "950" in short_soc or "Ascend950" in short_soc
		163	+ l0c_dtype = dtype if is_950 else "float32"
		164	+
		165	+ matrix_k = 16
		166	+ matrix_m = l0a_size // (matrix_k * dtype_bytes)
		167	+ matrix_n = l0b_size // (matrix_k * dtype_bytes)
		168	+
		169	+ input_bytes = 128 * 1024
		170	+ input_gm = tik_instance.Tensor(dtype, (input_bytes // dtype_bytes,), name="input_gm", scope=tik.scope_gm)
		171	+
		172	+ src_a1 = tik_instance.Tensor(dtype, (matrix_m, matrix_k), name="src_a1", scope=tik.scope_cbuf)
		173	+ src_b1 = tik_instance.Tensor(dtype, (matrix_k, matrix_k), name="src_b1", scope=tik.scope_cbuf)
		174	+ src_a2 = tik_instance.Tensor(dtype, (matrix_k, matrix_k), name="src_a2", scope=tik.scope_cbuf)
		175	+ src_b2 = tik_instance.Tensor(dtype, (matrix_k, matrix_n), name="src_b2", scope=tik.scope_cbuf)
		176	+ dst_l0c = tik_instance.Tensor(l0c_dtype, (matrix_m, matrix_k), name="dst_l0c", scope=tik.scope_cc)
		177	+
		178	+ a1_burst = matrix_m * matrix_k * dtype_bytes // 32
		179	+ b2_burst = matrix_k * matrix_n * dtype_bytes // 32
		180	+ small_burst = matrix_k * matrix_k * dtype_bytes // 32
		181	+
		182	+ with tik_instance.for_range(0, core_num, block_num=core_num):
		183	+ tik_instance.data_move(dst=src_a1, src=input_gm, sid=0, nburst=1, burst=a1_burst, src_stride=0, dst_stride=0)
		184	+ tik_instance.data_move(dst=src_b1, src=input_gm, sid=0, nburst=1, burst=small_burst, src_stride=0, dst_stride=0)
		185	+ tik_instance.matmul(dst=dst_l0c, a=src_a1, b=src_b1, m=matrix_m, k=matrix_k, n=matrix_k, init_l1out=True)
		186	+
		187	+ tik_instance.data_move(dst=src_a2, src=input_gm, sid=0, nburst=1, burst=small_burst, src_stride=0, dst_stride=0)
		188	+ tik_instance.data_move(dst=src_b2, src=input_gm, sid=0, nburst=1, burst=b2_burst, src_stride=0, dst_stride=0)
		189	+ tik_instance.matmul(dst=dst_l0c, a=src_a2, b=src_b2, m=matrix_k, k=matrix_k, n=matrix_n, init_l1out=True)
		190	+
		191	+ try:
		192	+ tik_instance.BuildCCE(kernel_name=kernel_name, inputs=(input_gm,), outputs=())
		193	+ finally:
		194	+ _pi.api_check_support = _orig_check
		195	+ _matmul_mod.api_check_support = _orig_check
		196	+ _load2d_mod.api_check_support = _orig_check
		197	+
		198	+ return tik_instance
		199	+
		200	+
126	def warmup(full_soc_version: str, kernel_name: str = "warmup"):	201	def warmup(full_soc_version: str, kernel_name: str = "warmup"):
127	"""	202	"""
128	a kernel without any action, but warmup all the cores, like: smmu/biu/tlb,	203	a kernel without any action, but warmup all the cores, like: smmu/biu/tlb,

Mttk/core_modules/operator/tbe_interface.py+123-5

             "api_config",
         ):
             return getattr(self._opc, item)
++        return super().__getattribute__(item)
+-            return super().__getattribute__(item)
     @property
     def core_type(self) -> str:
         self._all_opc_invoke("set_compile_soc_info", dev_plat, self._core_type)
     def switch_opc(self, opc_type: str):
++        if opc_type not in self.OpcImplement:
             raise ValueError(f"Invalid opc type: {opc_type}")
         self._opc = getattr(self, f"_{opc_type}_opc")
                     f"Compile clear_l1 failed. kernel or json file does not exist in {switches.kernel_meta}"
                 )
++    def compile_l0_clear(self):
++        switches = get_global_storage()
++        opc = self._tbe_opc
++        if switches.force_clear_l0 is not None:
++            obj_file = os.path.join(switches.kernel_meta, "clear_l0.o")
++            if os.path.exists(obj_file):
++                os.remove(obj_file)
++            full_soc_version = opc.get_soc_spec("FULL_SOC_VERSION")
++            short_soc_version = opc.get_soc_spec("SHORT_SOC_VERSION")
++
++            from ttk.utilities.platform import PLATFORM_BEFORE_DAVID
++
++            if short_soc_version not in PLATFORM_BEFORE_DAVID:
++                self._compile_l0_clear_ascendc(switches.kernel_meta, full_soc_version)
++            else:
++                from .helper_kernels import clear_l0
++
++                with opc.op_context.OpContext("pre-static") as cxt:
++                    tensor = {"shape": (1,), "range": ((1, None),), "dtype": "float16", "format": "ND"}
++                    attrs = {"full_soc_version": full_soc_version, "kernel_name": "clear_l0"}
++                    op_info = opc.op_info.OpInfo("ClearL0", "ClearL0")
++                    cxt.add_op_info(op_info)
++                    cxt.add_addition("op_name", "ClearL0")
++                    clear_l0(tensor, **attrs)
++            if not os.path.exists(obj_file) or not os.path.exists(os.path.join(switches.kernel_meta, "clear_l0.json")):
++                raise RuntimeError(
++                    f"Compile clear_l0 failed. kernel or json file does not exist in {switches.kernel_meta}"
++                )
++
++    @staticmethod
++    def _compile_l0_clear_ascendc(kernel_meta: str, full_soc_version: str):
++        """
++        David-generation chips have no TIK cube intrinsics (mmad/load2dv2/fixpipe),
++        so the clear_l0 helper kernel is implemented with Ascend C tensor_api and
++        compiled with ccec directly (see ascendc_kernels/clear_l0_ascendc.cpp).
++        Tile dimensions and --npu-arch are derived from the target chip's platform
++        config so the kernel adapts to different L0A/L0B/L0C sizes.
++        """
++        import json
++        import subprocess
++
++        import tbe.common.platform as tbe_platform
++
++        from ttk.utilities.platform import get_l0_clear_tile_params
++
++        tbe_platform.set_current_compile_soc_info(full_soc_version, "AiCore")
++        core_num = tbe_platform.get_soc_spec("CORE_NUM")
++        m_dim, k_dim, n_dim, npu_arch = get_l0_clear_tile_params(full_soc_version)
++
++        ascend_home = os.environ.get("ASCEND_HOME_PATH", "")
++        asc_dir = os.path.join(ascend_home, "asc")
++        os.makedirs(kernel_meta, mode=0o700, exist_ok=True)
++        src_file = os.path.join(os.path.dirname(__file__), "ascendc_kernels", "clear_l0_ascendc.cpp")
++        i_obj_file = os.path.join(kernel_meta, "clear_l0.i")
++
++        ccec_cmd = [
++            "ccec",
++            "-O2",
++            "--asc-aicore-lang",
++            src_file,
++            f"-DM_DIM={m_dim}",
++            f"-DK_DIM={k_dim}",
++            f"-DN_DIM={n_dim}",
++            "-I",
++            asc_dir,
++            "-I",
++            os.path.join(asc_dir, "include"),
++            "-I",
++            os.path.join(asc_dir, "include", "basic_api"),
++            f"--npu-arch={npu_arch}",
++            "--cce-aicore-only",
++            "-o",
++            i_obj_file,
++            "-mllvm",
++            "--cce-aicore-jump-expand=true",
++            "-mllvm",
++            "-cce-aicore-addr-transform",
++            "-mllvm",
++            "-cce-aicore-stack-size=0x8000",
++            "-mllvm",
++            "-cce-aicore-function-stack-size=0x8000",
++            "-mllvm",
++            "-cce-aicore-record-overflow=false",
++            "-mllvm",
++            "-cce-aicore-dcci-insert-for-scalar=false",
++            "-mllvm",
++            "-cce-aicore-dcci-before-kernel-end=false",
++        ]
++        result = subprocess.run(ccec_cmd, capture_output=True, text=True, check=False)
++        if result.returncode != 0:
++            raise RuntimeError(f"clear_l0 ccec compilation failed: {result.stderr}")
++
++        ld_cmd = [
++            "ld.lld",
++            "-m",
++            "aicorelinux",
++            "-Ttext=0",
++            i_obj_file,
++            "-static",
++            "-o",
++            os.path.join(kernel_meta, "clear_l0.o"),
++        ]
++        result = subprocess.run(ld_cmd, capture_output=True, text=True, check=False)
++        if result.returncode != 0:
++            raise RuntimeError(f"clear_l0 ld.lld failed: {result.stderr}")
++        if os.path.exists(i_obj_file):
++            os.remove(i_obj_file)
++
++        kernel_info = {
++            "binFileName": "clear_l0",
++            "binFileSuffix": ".o",
++            "blockDim": core_num,
++            "coreType": "AiCore",
++            "kernelName": "clear_l0",
++            "magic": "RT_DEV_BINARY_MAGIC_ELF",
++        }
++        with open(os.path.join(kernel_meta, "clear_l0.json"), "w", encoding="UTF-8") as f:
++            json.dump(kernel_info, f, indent=4)
++
     def compile_warmup_kernel(self):
         switches = get_global_storage()
         opc = self._tbe_opc
                 )
     def _get_any_opc(self) -> Optional[IOpc]:
++        for k in self.OpcImplement:
             opc = getattr(self, f"_{k}_opc")
             if opc.is_initialized():
                 return opc
         return None
     def _all_opc_invoke(self, func, *args, **kwargs):
++        for k in self.OpcImplement:
             opc: IOpc = getattr(self, f"_{k}_opc")
             if not opc.is_initialized():
                 continue

Mttk/core_modules/runtime/rts_interface.py+72-35

	@@ -530,23 +530,22 @@ class RTSInterface(RTSInterfaceBase):
530	ctypes.c_void_p(ctypes.addressof(c_device_ids)),	530	ctypes.c_void_p(ctypes.addressof(c_device_ids)),
531	)	531	)
532	self.prof_switch_version = 0	532	self.prof_switch_version = 0
		533	+ elif self.prof_switch_version == 0:
		534	+ self.api_call(
		535	+ "rtProfilerStart",
		536	+ None,
		537	+ c_prof_config,
		538	+ ctypes.c_int32(1),
		539	+ ctypes.c_void_p(ctypes.addressof(c_device_ids)),
		540	+ )
		541	+ elif self.prof_switch_version == 1:
		542	+ self.set_prof_switch(
		543	+ rts_info.MsprofCommandHandleType.PROF_COMMANDHANDLE_TYPE_START, c_device_ids, c_prof_config
		544	+ )
533	else:	545	else:
534	- if self.prof_switch_version == 0:	546	+ self.set_prof_switch_v2(
535	- self.api_call(	547	+ rts_info.MsprofCommandHandleType.PROF_COMMANDHANDLE_TYPE_START, c_device_ids, c_prof_config
536	- "rtProfilerStart",	548	+ )
537	- None,
538	- c_prof_config,
539	- ctypes.c_int32(1),
540	- ctypes.c_void_p(ctypes.addressof(c_device_ids)),
541	- )
542	- elif self.prof_switch_version == 1:
543	- self.set_prof_switch(
544	- rts_info.MsprofCommandHandleType.PROF_COMMANDHANDLE_TYPE_START, c_device_ids, c_prof_config
545	- )
546	- else:
547	- self.set_prof_switch_v2(
548	- rts_info.MsprofCommandHandleType.PROF_COMMANDHANDLE_TYPE_START, c_device_ids, c_prof_config
549	- )
550		549
551	def set_prof_switch(self, prof_switch_command_type, c_device_ids, c_prof_config):	550	def set_prof_switch(self, prof_switch_command_type, c_device_ids, c_prof_config):
552	command_type = prof_switch_command_type.value	551	command_type = prof_switch_command_type.value
	@@ -651,6 +650,32 @@ class RTSInterface(RTSInterfaceBase):
651	input_np_array = numpy.array([clean_val.item(0)] * (128 * 1024 // dtype_bytes), dtype=clean_val.dtype)	650	input_np_array = numpy.array([clean_val.item(0)] * (128 * 1024 // dtype_bytes), dtype=clean_val.dtype)
652	self._launch_helper_kernel(os.path.join(switches.kernel_meta, "clear_l1.o"), args=(input_np_array,))	651	self._launch_helper_kernel(os.path.join(switches.kernel_meta, "clear_l1.o"), args=(input_np_array,))
653		652
		653	+ def clear_l0(self, switches: "ttk.utilities.SWITCHES"):
		654	+ if switches.force_clear_l0 is None or switches.mode.is_model():
		655	+ return
		656	+ clean_val = switches.force_clear_l0
		657	+ verify_arg = None
		658	+ from ttk.utilities.platform import PLATFORM_BEFORE_DAVID, get_l0_clear_tile_params
		659	+
		660	+ if self.short_soc_version not in PLATFORM_BEFORE_DAVID:
		661	+ m_dim, k_dim, n_dim, _ = get_l0_clear_tile_params(switches.dev_plat)
		662	+ input_np_array = numpy.array(
		663	+ [numpy.float16(clean_val.item(0))] * (m_dim * k_dim + k_dim * n_dim), dtype=numpy.float16
		664	+ )
		665	+ # David-generation chips use an Ascend C tensor_api kernel
		666	+ # (two GM args: input and fixpipe output)
		667	+ output_np_array = numpy.zeros(m_dim * n_dim, dtype=numpy.float32)
		668	+ args = (input_np_array, output_np_array)
		669	+ # fp16 x fp16 -> fp32 mmad of all-clean_val matrices: every element == K * v^2 exactly
		670	+ v = numpy.float32(numpy.float16(clean_val.item(0)))
		671	+ verify_arg = (1, v * v * numpy.float32(k_dim))
		672	+ else:
		673	+ input_np_array = numpy.array([numpy.float16(clean_val.item(0))] * (128 * 1024 // 2), dtype=numpy.float16)
		674	+ args = (input_np_array,)
		675	+ logging.info(f"[clear_l0] launching helper kernel, clean_val={clean_val}, device={self.device_id}")
		676	+ self._launch_helper_kernel(os.path.join(switches.kernel_meta, "clear_l0.o"), args=args, verify_arg=verify_arg)
		677	+ logging.info("[clear_l0] helper kernel launched successfully")
		678	+
654	def clear_ub(self, switches: "ttk.utilities.SWITCHES"):	679	def clear_ub(self, switches: "ttk.utilities.SWITCHES"):
655	if switches.force_clear_ub is None or switches.mode.is_model():	680	if switches.force_clear_ub is None or switches.mode.is_model():
656	return	681	return
	@@ -674,6 +699,7 @@ class RTSInterface(RTSInterfaceBase):
674	kernel: str,	699	kernel: str,
675	args: Optional[Union[type(None), numpy.ndarray]] = None,	700	args: Optional[Union[type(None), numpy.ndarray]] = None,
676	print_arg_indices: Optional[tuple] = None,	701	print_arg_indices: Optional[tuple] = None,
		702	+ verify_arg: Optional[tuple] = None,
677	):	703	):
678	stream = self.create_stream()	704	stream = self.create_stream()
679	kernel_name = Path(kernel).stem	705	kernel_name = Path(kernel).stem
	@@ -706,6 +732,22 @@ class RTSInterface(RTSInterfaceBase):
706	byte_size = int(math.ceil(np_array.size * get_dtype_width(np_array.dtype)))	732	byte_size = int(math.ceil(np_array.size * get_dtype_width(np_array.dtype)))
707	byte_array = self.get_data_from_hbm(npu_ptr, byte_size)	733	byte_array = self.get_data_from_hbm(npu_ptr, byte_size)
708	print(numpy.frombuffer(byte_array, dtype=np_array.dtype))	734	print(numpy.frombuffer(byte_array, dtype=np_array.dtype))
		735	+ if verify_arg is not None:
		736	+ idx, expected = verify_arg
		737	+ npu_ptr, np_array = dev_mem_addrs[idx], args[idx]
		738	+ byte_size = int(math.ceil(np_array.size * get_dtype_width(np_array.dtype)))
		739	+ got = numpy.frombuffer(self.get_data_from_hbm(npu_ptr, byte_size), dtype=np_array.dtype)
		740	+ mismatch = int((got != expected).sum())
		741	+ if mismatch == 0:
		742	+ logging.info(
		743	+ f"[{kernel_name}] device result verified: all {got.size} elems == {expected}, "
		744	+ "full GM->L1->L0A/L0B->Mmad->L0C->fixpipe chain executed"
		745	+ )
		746	+ else:
		747	+ logging.warning(
		748	+ f"[{kernel_name}] device result VERIFICATION FAILED: {mismatch}/{got.size} "
		749	+ f"elems != {expected}"
		750	+ )
709	except Exception as e:	751	except Exception as e:
710	logging.exception(f"synchronize_with_stream [{kernel_name}] failed. {e}")	752	logging.exception(f"synchronize_with_stream [{kernel_name}] failed. {e}")
711	finally:	753	finally:
	@@ -799,8 +841,7 @@ class RTSInterface(RTSInterfaceBase):
799	self.api_call("rtGetC2cCtrlAddr", None, ret_addr, ret_len)	841	self.api_call("rtGetC2cCtrlAddr", None, ret_addr, ret_len)
800	ffts_addr = ctypes.c_void_p(ret_addr[0])	842	ffts_addr = ctypes.c_void_p(ret_addr[0])
801	return 0 if self.is_model() and ffts_addr.value is None else ffts_addr	843	return 0 if self.is_model() and ffts_addr.value is None else ffts_addr
802	- else:	844	+ return None
803	- return None
804		845
805	def _launch_kernel_v2(self, launch_args: rts_structures.LaunchKernelArgs, stream: Optional[ctypes.c_void_p] = None):	846	def _launch_kernel_v2(self, launch_args: rts_structures.LaunchKernelArgs, stream: Optional[ctypes.c_void_p] = None):
806	launch_args.insert_ffts_addr(self._get_c2c_ctrl_addr(launch_args.mix_kernel))	847	launch_args.insert_ffts_addr(self._get_c2c_ctrl_addr(launch_args.mix_kernel))
	@@ -864,10 +905,9 @@ class RTSInterface(RTSInterfaceBase):
864	def _build_rt_task_cfg_info(self, simt_share_memory_size: int, schedule_mode: int):	905	def _build_rt_task_cfg_info(self, simt_share_memory_size: int, schedule_mode: int):
865	if simt_share_memory_size <= 0 and schedule_mode == 0:	906	if simt_share_memory_size <= 0 and schedule_mode == 0:
866	return None	907	return None
867	- elif self._is_0903_branch():	908	+ if self._is_0903_branch():
868	return rts_structures.RtTaskCfgInfoBranch0903(simt_share_memory_size, schedule_mode)	909	return rts_structures.RtTaskCfgInfoBranch0903(simt_share_memory_size, schedule_mode)
869	- else:	910	+ return rts_structures.RtTaskCfgInfo(simt_share_memory_size, schedule_mode)
870	- return rts_structures.RtTaskCfgInfo(simt_share_memory_size, schedule_mode)
871		911
872	def _set_simt_stack_size(self, typ: int, stack_size: int, device_id: Optional[int] = None):	912	def _set_simt_stack_size(self, typ: int, stack_size: int, device_id: Optional[int] = None):
873	if stack_size < 0:	913	if stack_size < 0:
	@@ -895,19 +935,19 @@ class RTSInterface(RTSInterfaceBase):
895	if "AICORE_TRAP_EXCEPTION" in exception:	935	if "AICORE_TRAP_EXCEPTION" in exception:
896	logging.error("Reached AICORE Trap Exception")	936	logging.error("Reached AICORE Trap Exception")
897	return "TRAP"	937	return "TRAP"
898	- elif "AICORE_EXCEPTION" in exception:	938	+ if "AICORE_EXCEPTION" in exception:
899	logging.error("AIC_ERROR encountered")	939	logging.error("AIC_ERROR encountered")
900	return "AIC_ERROR"	940	return "AIC_ERROR"
901	- elif "VECTOR_CORE_EXCEPTION" in exception:	941	+ if "VECTOR_CORE_EXCEPTION" in exception:
902	logging.error("VEC_ERROR encountered")	942	logging.error("VEC_ERROR encountered")
903	return "VEC_ERROR"	943	return "VEC_ERROR"
904	- elif "AICORE_TIMEOUT" in exception or "RT_STREAM_SYNC_TIMEOUT" in exception:	944	+ if "AICORE_TIMEOUT" in exception or "RT_STREAM_SYNC_TIMEOUT" in exception:
905	logging.error("AIC Task TIMEOUT")	945	logging.error("AIC Task TIMEOUT")
906	return "TIMEOUT"	946	return "TIMEOUT"
907	- elif "HEARTBEAT" in exception:	947	+ if "HEARTBEAT" in exception:
908	logging.critical("Detected critical device heartbeat lost exception, process will halt.")	948	logging.critical("Detected critical device heartbeat lost exception, process will halt.")
909	if shutil.which("msnpureport") is not None:	949	if shutil.which("msnpureport") is not None:
910	- os.system(f"mkdir -p errors/{os.getpid()} && cd errors/{os.getpid()} && msnpureport && cd -")	950	+ os.system(f"mkdir -p errors/{os.getpid()} && cd errors/{os.getpid()} && msnpureport && cd -") # noqa: S605
911	while True:	951	while True:
912	time.sleep(10)	952	time.sleep(10)
913	logging.critical("This testcase is killing device!!!! AND DEVICE WAS ALREADY DEAD")	953	logging.critical("This testcase is killing device!!!! AND DEVICE WAS ALREADY DEAD")
	@@ -921,8 +961,7 @@ class RTSInterface(RTSInterfaceBase):
921	raise TypeError(f"Copy numpy array to hbm supports ndarray only, but received {str(type(_nparray))}")	961	raise TypeError(f"Copy numpy array to hbm supports ndarray only, but received {str(type(_nparray))}")
922	if is_4bit_dtype(_nparray.dtype):	962	if is_4bit_dtype(_nparray.dtype):
923	return pack_4bits(_nparray)	963	return pack_4bits(_nparray)
924	- else:	964	+ return _nparray.flatten()
925	- return _nparray.flatten()
926		965
927	@staticmethod	966	@staticmethod
928	def _parse_error_code(error_type: int, error_code: int) -> str:	967	def _parse_error_code(error_type: int, error_code: int) -> str:
	@@ -981,18 +1020,16 @@ class RTSInterface(RTSInterfaceBase):
981	def int_magic(magic: str) -> int:	1020	def int_magic(magic: str) -> int:
982	if magic in rts_info.rt_binary_magic_dict:	1021	if magic in rts_info.rt_binary_magic_dict:
983	return rts_info.rt_binary_magic_dict[magic]	1022	return rts_info.rt_binary_magic_dict[magic]
984	- else:	1023	+ raise RuntimeError(f"Unknown kernel magic: {magic}")
985	- raise RuntimeError(f"Unknown kernel magic: {magic}")
986		1024
987	@staticmethod	1025	@staticmethod
988	def core_type_to_magic(core_type: str) -> int:	1026	def core_type_to_magic(core_type: str) -> int:
989	if core_type == "AiCore":	1027	if core_type == "AiCore":
990	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF"]	1028	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF"]
991	- elif core_type == "VectorCore":	1029	+ if core_type == "VectorCore":
992	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF_AIVEC"]	1030	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF_AIVEC"]
993	- elif core_type == "CubeCore":	1031	+ if core_type == "CubeCore":
994	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF_AICUBE"]	1032	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF_AICUBE"]
995	- elif core_type == "AiCpu":	1033	+ if core_type == "AiCpu":
996	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF_AICPU"]	1034	return rts_info.rt_binary_magic_dict["RT_DEV_BINARY_MAGIC_ELF_AICPU"]
997	- else:	1035	+ raise RuntimeError(f"Unknown core type: {core_type}")
998	- raise RuntimeError(f"Unknown core type: {core_type}")

Mttk/utilities/classes.py+15-16

	@@ -44,7 +44,7 @@ class MODE(Enum):
44	return None	44	return None
45		45
46	def is_online_board(self) -> bool:	46	def is_online_board(self) -> bool:
47	- return True if self in [MODE.ASCEND_ONBOARD] else False	47	+ return self in [MODE.ASCEND_ONBOARD]
48		48
49	def has_device(self) -> bool:	49	def has_device(self) -> bool:
50	return self.is_online_board()	50	return self.is_online_board()
	@@ -145,6 +145,7 @@ class SWITCHES:
145	"force_block_dim",	145	"force_block_dim",
146	"force_clear_ub",	146	"force_clear_ub",
147	"force_clear_l1",	147	"force_clear_l1",
		148	+ "force_clear_l0",
148	"force_simt_ub_size",	149	"force_simt_ub_size",
149	"progress_output",	150	"progress_output",
150	"proc_no_reuse",	151	"proc_no_reuse",
	@@ -214,8 +215,7 @@ class SWITCHES:
214	def run_time(self) -> int:	215	def run_time(self) -> int:
215	if self.mode.is_model():	216	if self.mode.is_model():
216	return self._run_time or 1	217	return self._run_time or 1
217	- else:	218	+ return self._run_time or 3
218	- return self._run_time or 3
219		219
220	@run_time.setter	220	@run_time.setter
221	def run_time(self, val):	221	def run_time(self, val):
	@@ -276,6 +276,7 @@ class SWITCHES:
276	self.force_block_dim = [None, None, None]	276	self.force_block_dim = [None, None, None]
277	self.force_clear_ub = None	277	self.force_clear_ub = None
278	self.force_clear_l1 = None	278	self.force_clear_l1 = None
		279	+ self.force_clear_l0 = None
279	self.force_simt_ub_size = [None, None, None] # dynamic/const/binary	280	self.force_simt_ub_size = [None, None, None] # dynamic/const/binary
280	self.proc_no_reuse = False	281	self.proc_no_reuse = False
281	# Hidden switches	282	# Hidden switches
	@@ -368,13 +369,13 @@ class SubKernelJsonInfo:
368		369
369	@classmethod	370	@classmethod
370	def from_dict(cls, json_dict: dict):	371	def from_dict(cls, json_dict: dict):
371	- parameters = json_dict.get("parameters", None)	372	+ parameters = json_dict.get("parameters")
372	if parameters is not None:	373	if parameters is not None:
373	parameters = tuple(parameters)	374	parameters = tuple(parameters)
374	- magic = json_dict.get("magic", None)	375	+ magic = json_dict.get("magic")
375	- core_type = json_dict.get("coreType", None)	376	+ core_type = json_dict.get("coreType")
376	kernel_name = json_dict["kernelName"]	377	kernel_name = json_dict["kernelName"]
377	- task_ration = json_dict.get("taskRation", None)	378	+ task_ration = json_dict.get("taskRation")
378	if task_ration is not None:	379	if task_ration is not None:
379	if not isinstance(task_ration, str) or ":" not in task_ration:	380	if not isinstance(task_ration, str) or ":" not in task_ration:
380	raise ValueError(f"task_ration [{task_ration}] is invalid. It may be a bug of compiler.")	381	raise ValueError(f"task_ration [{task_ration}] is invalid. It may be a bug of compiler.")
	@@ -517,14 +518,13 @@ class KernelJsonInfo:
517	sub_kernel_name = f"{self.kernel_name}_{tiling_key}"	518	sub_kernel_name = f"{self.kernel_name}_{tiling_key}"
518	if sub_kernel_name not in self.sub_kernels:	519	if sub_kernel_name not in self.sub_kernels:
519	return self	520	return self
520	- elif sub_kernel_name in self._sub_kernel_cache:	521	+ if sub_kernel_name in self._sub_kernel_cache:
521	return self._sub_kernel_cache[sub_kernel_name]	522	return self._sub_kernel_cache[sub_kernel_name]
522	- else:	523	+ sk = self.sub_kernels[sub_kernel_name]
523	- sk = self.sub_kernels[sub_kernel_name]	524	+ ret = copy.deepcopy(self)
524	- ret = copy.deepcopy(self)	525	+ self._migrate(sk, ret)
525	- self._migrate(sk, ret)	526	+ self._sub_kernel_cache[sub_kernel_name] = ret
526	- self._sub_kernel_cache[sub_kernel_name] = ret	527	+ return ret
527	- return ret
528		528
529	def dynamic_param_is_folded(self):	529	def dynamic_param_is_folded(self):
530	return self.dynamic_param_mode == "folded_with_desc"	530	return self.dynamic_param_mode == "folded_with_desc"
	@@ -725,8 +725,7 @@ class DynamicCompilationResult(BaseCompilationResult):
725	def block_dim(self) -> int:	725	def block_dim(self) -> int:
726	if not self.tiling_result or not self.tiling_result.block_dim:	726	if not self.tiling_result or not self.tiling_result.block_dim:
727	return 0	727	return 0
728	- else:	728	+ return self.tiling_result.block_dim
729	- return self.tiling_result.block_dim
730		729
731	@block_dim.setter	730	@block_dim.setter
732	def block_dim(self, value):	731	def block_dim(self, value):

Mttk/utilities/platform.py+93-17

	@@ -21,6 +21,7 @@ __all__ = [
21	"get_ascend_full_soc_version",	21	"get_ascend_full_soc_version",
22	"get_npu_hw_info",	22	"get_npu_hw_info",
23	"get_npu_available_device_ids",	23	"get_npu_available_device_ids",
		24	+ "get_l0_clear_tile_params",
24	"PLATFORM_BEFORE_DAVID",	25	"PLATFORM_BEFORE_DAVID",
25	]	26	]
26		27
	@@ -63,7 +64,7 @@ def get_opp_paths(source: str) -> list:
63	if not opp_path:	64	if not opp_path:
64	raise RuntimeError("Environment `ASCEND_OPP_PATH` is not set.")	65	raise RuntimeError("Environment `ASCEND_OPP_PATH` is not set.")
65	return [opp_path]	66	return [opp_path]
66	- elif source == "vendor":	67	+ if source == "vendor":
67	opp_path = os.getenv("ASCEND_OPP_PATH", "")	68	opp_path = os.getenv("ASCEND_OPP_PATH", "")
68	if not opp_path:	69	if not opp_path:
69	return []	70	return []
	@@ -74,12 +75,12 @@ def get_opp_paths(source: str) -> list:
74	return [custom_path] if os.path.isdir(custom_path) else []	75	return [custom_path] if os.path.isdir(custom_path) else []
75	vendors = []	76	vendors = []
76	with open(config_file) as f:	77	with open(config_file) as f:
77	- for line in f:	78	+ for raw_line in f:
78	- line = line.strip()	79	+ stripped_line = raw_line.strip()
79	- if not line.startswith("load_priority="):	80	+ if not stripped_line.startswith("load_priority="):
80	continue	81	continue
81	- for name in line.split("=", 1)[1].split(","):	82	+ for raw_name in stripped_line.split("=", 1)[1].split(","):
82	- name = name.strip()	83	+ name = raw_name.strip()
83	if not name:	84	if not name:
84	continue	85	continue
85	p = os.path.join(opp_path, "vendors", name)	86	p = os.path.join(opp_path, "vendors", name)
	@@ -87,7 +88,7 @@ def get_opp_paths(source: str) -> list:
87	vendors.append(p)	88	vendors.append(p)
88	break	89	break
89	return vendors	90	return vendors
90	- elif source == "custom":	91	+ if source == "custom":
91	env_val = os.getenv("ASCEND_CUSTOM_OPP_PATH", "")	92	env_val = os.getenv("ASCEND_CUSTOM_OPP_PATH", "")
92	if not env_val:	93	if not env_val:
93	return []	94	return []
	@@ -110,9 +111,9 @@ def get_op_impl_paths(source: str) -> list:
110	"""Return {opp}/op_impl paths. source: 'builtin' \| 'vendor' \| 'custom'. Always returns list."""	111	"""Return {opp}/op_impl paths. source: 'builtin' \| 'vendor' \| 'custom'. Always returns list."""
111	if source == "builtin":	112	if source == "builtin":
112	return [os.path.join(get_opp_paths("builtin")[0], _builtin_op_impl_rel())]	113	return [os.path.join(get_opp_paths("builtin")[0], _builtin_op_impl_rel())]
113	- elif source == "vendor":	114	+ if source == "vendor":
114	return [os.path.join(p, "op_impl") for p in get_opp_paths("vendor")]	115	return [os.path.join(p, "op_impl") for p in get_opp_paths("vendor")]
115	- elif source == "custom":	116	+ if source == "custom":
116	return [os.path.join(p, "op_impl") for p in get_opp_paths("custom")]	117	return [os.path.join(p, "op_impl") for p in get_opp_paths("custom")]
117	raise ValueError(f"Unknown source: {source}, must be one of {_VALID_SOURCES}")	118	raise ValueError(f"Unknown source: {source}, must be one of {_VALID_SOURCES}")
118		119
	@@ -122,9 +123,9 @@ def get_impl_base_paths(source: str) -> list:
122	"""Return {opp}/op_impl/ai_core/tbe paths. source: 'builtin' \| 'vendor' \| 'custom'. Always returns list."""	123	"""Return {opp}/op_impl/ai_core/tbe paths. source: 'builtin' \| 'vendor' \| 'custom'. Always returns list."""
123	if source == "builtin":	124	if source == "builtin":
124	return [os.path.join(get_opp_paths("builtin")[0], _builtin_op_impl_rel(), "ai_core", "tbe")]	125	return [os.path.join(get_opp_paths("builtin")[0], _builtin_op_impl_rel(), "ai_core", "tbe")]
125	- elif source == "vendor":	126	+ if source == "vendor":
126	return [os.path.join(p, "op_impl", "ai_core", "tbe") for p in get_opp_paths("vendor")]	127	return [os.path.join(p, "op_impl", "ai_core", "tbe") for p in get_opp_paths("vendor")]
127	- elif source == "custom":	128	+ if source == "custom":
128	return [os.path.join(p, "op_impl", "ai_core", "tbe") for p in get_opp_paths("custom")]	129	return [os.path.join(p, "op_impl", "ai_core", "tbe") for p in get_opp_paths("custom")]
129	raise ValueError(f"Unknown source: {source}, must be one of {_VALID_SOURCES}")	130	raise ValueError(f"Unknown source: {source}, must be one of {_VALID_SOURCES}")
130		131
	@@ -136,7 +137,7 @@ def get_ascend_scene_info() -> Tuple[str, str]:
136	scene_os, scene_arch = "", ""	137	scene_os, scene_arch = "", ""
137	if os.path.isfile(scene_file):	138	if os.path.isfile(scene_file):
138	with open(scene_file) as f:	139	with open(scene_file) as f:
139	- scene_info = list(map(lambda x: x.strip(), f.readlines()))	140	+ scene_info = [line.strip() for line in f.readlines()]
140	for item_info in scene_info:	141	for item_info in scene_info:
141	if "os=" in item_info:	142	if "os=" in item_info:
142	scene_os = item_info.split("=")[-1]	143	scene_os = item_info.split("=")[-1]
	@@ -150,9 +151,9 @@ def get_op_info_paths(source: str, soc_lower: str) -> list:
150	"""Return {impl_base}/config/{soc_lower} paths. source: 'builtin' \| 'vendor' \| 'custom'. Always returns list."""	151	"""Return {impl_base}/config/{soc_lower} paths. source: 'builtin' \| 'vendor' \| 'custom'. Always returns list."""
151	if source == "builtin":	152	if source == "builtin":
152	return [os.path.join(get_impl_base_paths("builtin")[0], "config", soc_lower)]	153	return [os.path.join(get_impl_base_paths("builtin")[0], "config", soc_lower)]
153	- elif source == "vendor":	154	+ if source == "vendor":
154	return [os.path.join(p, "op_impl", "ai_core", "tbe", "config", soc_lower) for p in get_opp_paths("vendor")]	155	return [os.path.join(p, "op_impl", "ai_core", "tbe", "config", soc_lower) for p in get_opp_paths("vendor")]
155	- elif source == "custom":	156	+ if source == "custom":
156	return [os.path.join(p, "op_impl", "ai_core", "tbe", "config", soc_lower) for p in get_opp_paths("custom")]	157	return [os.path.join(p, "op_impl", "ai_core", "tbe", "config", soc_lower) for p in get_opp_paths("custom")]
157	raise ValueError(f"Unknown source: {source}, must be one of {_VALID_SOURCES}")	158	raise ValueError(f"Unknown source: {source}, must be one of {_VALID_SOURCES}")
158		159
	@@ -192,7 +193,8 @@ def get_npu_hw_info(full_soc_version):
192	Returns:	193	Returns:
193	dict with keys: short_soc_version, ccec_aic_version, npu_arch,	194	dict with keys: short_soc_version, ccec_aic_version, npu_arch,
194	ai_core_cnt, cube_core_cnt, vector_core_cnt, cv_split,	195	ai_core_cnt, cube_core_cnt, vector_core_cnt, cv_split,
195	- core_type_list, support_bf16, support_inf_nan	196	+ core_type_list, support_bf16, support_inf_nan,
		197	+ l0_a_size, l0_b_size, l0_c_size
196		198
197	Raises:	199	Raises:
198	FileNotFoundError: if ini file not found	200	FileNotFoundError: if ini file not found
	@@ -213,6 +215,7 @@ def get_npu_hw_info(full_soc_version):
213	raw_fields = {	215	raw_fields = {
214	"short_soc_version": ("version", "Short_SoC_version"),	216	"short_soc_version": ("version", "Short_SoC_version"),
215	"ccec_aic_version": ("version", "CCEC_AIC_version"),	217	"ccec_aic_version": ("version", "CCEC_AIC_version"),
		218	+ "npu_arch": ("version", "NpuArch"),
216	"ai_core_cnt": ("SoCInfo", "ai_core_cnt"),	219	"ai_core_cnt": ("SoCInfo", "ai_core_cnt"),
217	"cube_core_cnt": ("SoCInfo", "cube_core_cnt"),	220	"cube_core_cnt": ("SoCInfo", "cube_core_cnt"),
218	"vector_core_cnt": ("SoCInfo", "vector_core_cnt"),	221	"vector_core_cnt": ("SoCInfo", "vector_core_cnt"),
	@@ -220,11 +223,14 @@ def get_npu_hw_info(full_soc_version):
220	"core_type_list": ("SoCInfo", "core_type_list"),	223	"core_type_list": ("SoCInfo", "core_type_list"),
221	"support_bf16": ("SoCInfo", "support_bf16"),	224	"support_bf16": ("SoCInfo", "support_bf16"),
222	"support_inf_nan": ("SoCInfo", "support_inf_nan"),	225	"support_inf_nan": ("SoCInfo", "support_inf_nan"),
		226	+ "l0_a_size": ("AICoreSpec", "l0_a_size"),
		227	+ "l0_b_size": ("AICoreSpec", "l0_b_size"),
		228	+ "l0_c_size": ("AICoreSpec", "l0_c_size"),
223	}	229	}
224		230
225	- int_fields = {"ai_core_cnt", "cube_core_cnt", "vector_core_cnt"}	231	+ int_fields = {"ai_core_cnt", "cube_core_cnt", "vector_core_cnt", "l0_a_size", "l0_b_size", "l0_c_size"}
226	bool_fields = {"support_bf16", "support_inf_nan"}	232	bool_fields = {"support_bf16", "support_inf_nan"}
227	- int_defaults = {"cube_core_cnt": 0, "vector_core_cnt": 0}	233	+ int_defaults = {"cube_core_cnt": 0, "vector_core_cnt": 0, "l0_a_size": 0, "l0_b_size": 0, "l0_c_size": 0}
228		234
229	result = {}	235	result = {}
230	missing = []	236	missing = []
	@@ -268,6 +274,76 @@ def get_npu_hw_info(full_soc_version):
268	return result	274	return result
269		275
270		276
		277	+@lru_cache(maxsize=None)
		278	+def get_l0_clear_tile_params(full_soc_version: str) -> Tuple[int, int, int, str]:
		279	+ """Compute M/K/N tile dimensions and npu_arch for the clear_l0 Ascend C helper kernel.
		280	+
		281	+ The kernel performs one mmad pass (fp16 A x fp16 B -> fp32 C) that covers the
		282	+ full L0A/L0B and as much of L0C as a single pass allows:
		283	+
		284	+ L0A tile: M * K * 2B (fp16, full L0A)
		285	+ L0B tile: K * N * 2B (fp16, full L0B)
		286	+ L0C tile: M * N * 4B (fp32, <= L0C)
		287	+
		288	+ K is chosen as the smallest multiple of 16 that:
		289	+ 1. Divides both L0A_SIZE / dtype_bytes and L0B_SIZE / dtype_bytes evenly.
		290	+ 2. Satisfies M * N * 4 <= L0C_SIZE.
		291	+
		292	+ Args:
		293	+ full_soc_version: e.g. 'Ascend950DT_9596'
		294	+
		295	+ Returns:
		296	+ (M_DIM, K_DIM, N_DIM, npu_arch) where npu_arch is the ccec --npu-arch value
		297	+ (e.g. 'dav-3510').
		298	+
		299	+ Raises:
		300	+ RuntimeError: if L0 sizes are zero or no valid K can be found.
		301	+ """
		302	+ import math
		303	+
		304	+ hw_info = get_npu_hw_info(full_soc_version)
		305	+ l0a_size = hw_info.get("l0_a_size", 0)
		306	+ l0b_size = hw_info.get("l0_b_size", 0)
		307	+ l0c_size = hw_info.get("l0_c_size", 0)
		308	+
		309	+ if l0a_size == 0 or l0b_size == 0 or l0c_size == 0:
		310	+ raise RuntimeError(
		311	+ f"[{full_soc_version}] L0 sizes are zero (l0a={l0a_size}, l0b={l0b_size}, l0c={l0c_size}), "
		312	+ "clear_l0 Ascend C kernel not supported"
		313	+ )
		314	+
		315	+ dtype_bytes = 2 # fp16
		316	+ l0c_dtype_bytes = 4 # fp32
		317	+
		318	+ l0a_elems = l0a_size // dtype_bytes
		319	+ l0b_elems = l0b_size // dtype_bytes
		320	+
		321	+ # Minimum K so that M * N * l0c_dtype_bytes <= l0c_size
		322	+ # M = l0a_elems / K, N = l0b_elems / K
		323	+ # => l0a_elems * l0b_elems * l0c_dtype_bytes / K^2 <= l0c_size
		324	+ # => K >= sqrt(l0a_elems * l0b_elems * l0c_dtype_bytes / l0c_size)
		325	+ min_k_sq = math.ceil(l0a_elems * l0b_elems * l0c_dtype_bytes / l0c_size)
		326	+ min_k = max(16, math.isqrt(min_k_sq))
		327	+ # Round up to multiple of 16 (cube instruction alignment)
		328	+ k_dim = ((min_k + 15) // 16) * 16
		329	+ # K must divide both element counts evenly
		330	+ while k_dim <= min(l0a_elems, l0b_elems):
		331	+ if l0a_elems % k_dim == 0 and l0b_elems % k_dim == 0:
		332	+ break
		333	+ k_dim += 16
		334	+ else:
		335	+ raise RuntimeError(
		336	+ f"[{full_soc_version}] cannot find K (multiple of 16, divisor of "
		337	+ f"l0a_elems={l0a_elems} and l0b_elems={l0b_elems}) satisfying L0C constraint"
		338	+ )
		339	+
		340	+ m_dim = l0a_elems // k_dim
		341	+ n_dim = l0b_elems // k_dim
		342	+ npu_arch = f"dav-{hw_info.get('npu_arch', '')}"
		343	+
		344	+ return m_dim, k_dim, n_dim, npu_arch
		345	+
		346	+
271	@lru_cache(maxsize=None)	347	@lru_cache(maxsize=None)
272	def get_npu_available_device_ids():	348	def get_npu_available_device_ids():
273	"""Get list of available Ascend device IDs.	349	"""Get list of available Ascend device IDs.