[inductor]add inductor test-pytorch-AtomGit

Atest/_inductor/__init__.py+3-0

@@ -0,0 +1,3 @@
	1	+#!/usr/bin/env python
	2	+# -- coding: utf-8 --
	3	+# Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.

Atest/_inductor/test_abs.py+26-0

@@ -0,0 +1,26 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestAbs(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.abs(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', [(1024, 32), (256, 8)])
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+ self.assertEqual(std_result, inductor_result, atol=1e-3, rtol=1e-3)
	22	+
	23	+instantiate_parametrized_tests(TestAbs)
	24	+
	25	+if __name__ == "__main__":
	26	+ run_tests()

Atest/_inductor/test_add.py+29-0

@@ -0,0 +1,29 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestAdd(TestUtils):
	8	+ def op_calc(self, first_element, second_element):
	9	+ result = first_element + second_element
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+ second_element = self._generate_tensor(shape, dtype)
	17	+
	18	+ std_sum = self.op_calc(first_element, second_element)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_sum = compiled_op_calc(first_element, second_element)
	22	+
	23	+ self.assertEqual(std_sum, inductor_sum)
	24	+
	25	+
	26	+instantiate_parametrized_tests(TestAdd)
	27	+
	28	+if __name__ == "__main__":
	29	+ run_tests()

Atest/_inductor/test_add_sum.py+38-0

@@ -0,0 +1,38 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSumAdd(TestUtils):
	8	+ def foo(self, a, b, dim):
	9	+ y = a + b
	10	+ y = y.sum(dim)
	11	+ return y
	12	+
	13	+ # case：change shapes
	14	+ @parametrize('shape', [(9, 9, 31, 64)])
	15	+ @parametrize('dim', [3])
	16	+ @parametrize('dtype', ['float32'])
	17	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	18	+ a, b = [torch.randn(shape, requires_grad=False, dtype=torch.float32, device="npu") for _ in range(2)]
	19	+ r1 = self.foo(a, b, dim)
	20	+ func = torch.compile(self.foo, backend="inductor", dynamic=False)
	21	+ r = func(a, b, dim)
	22	+ self.assertEqual(r, r1, atol=1e-3, rtol=1e-3)
	23	+
	24	+ @parametrize('shape', [(9, 10, 31, 63)])
	25	+ @parametrize('dim', [0, 1])
	26	+ @parametrize('dtype', ['float32'])
	27	+ def test_reduction_cases_shapes1(self, shape, dim, dtype):
	28	+ a, b = [torch.randn(shape, requires_grad=False, dtype=torch.float32, device="npu") for _ in range(2)]
	29	+ r1 = self.foo(a, b, dim)
	30	+ func = torch.compile(self.foo, backend="inductor", dynamic=False)
	31	+ r = func(a, b, dim)
	32	+ self.assertEqual(r, r1, atol=1e-3, rtol=1e-3)
	33	+
	34	+
	35	+instantiate_parametrized_tests(TestSumAdd)
	36	+
	37	+if __name__ == "__main__":
	38	+ run_tests()

Atest/_inductor/test_alias.py+27-0

@@ -0,0 +1,27 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestAlias(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ x = torch.ops.aten.alias(input_element)
	10	+ y = x + 1.0
	11	+ return y
	12	+
	13	+ # case：change shapes
	14	+ @parametrize('shape', [(32, 64)])
	15	+ @parametrize('dim', [0])
	16	+ @parametrize('dtype', ['float32'])
	17	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	18	+ input_element = self._generate_tensor(shape, dtype)
	19	+ std_ret = self.op_calc(input_element, dim)
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_ret = compiled_op_calc(input_element, dim)
	22	+ self.assertEqual(std_ret, inductor_ret, atol=1e-1, rtol=1e-1, equal_nan=True)
	23	+
	24	+instantiate_parametrized_tests(TestAlias)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_argmax.py+24-0

@@ -0,0 +1,24 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestArgmax(TestUtils):
	8	+ def argmax(self, a, dim):
	9	+ return torch.argmax(a, dim)
	10	+
	11	+ def test_argmax(self):
	12	+ shape = (512, 64)
	13	+ dim = -1
	14	+ a = torch.randn(shape, requires_grad=False, dtype=torch.float32, device='npu')
	15	+
	16	+ argmax_triton = torch.compile(self.argmax, backend="inductor", dynamic=False)
	17	+ r = self.argmax(a, dim)
	18	+ r1 = argmax_triton(a, dim)
	19	+ self.assertEqual(r, r1, atol=1e-3, rtol=1e-3)
	20	+
	21	+instantiate_parametrized_tests(TestArgmax)
	22	+
	23	+if __name__ == "__main__":
	24	+ run_tests()

Atest/_inductor/test_argmax_unalign.py+24-0

@@ -0,0 +1,24 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestMaxWithIndex(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return torch.argmax(input_element, dim)
	10	+
	11	+ @parametrize('shape', [(512, 64)]) # (513, 64), (514,33)
	12	+ @parametrize('dim', [-1])
	13	+ @parametrize('dtype', ['float32'])
	14	+ def test_reduction_cases(self, shape, dim, dtype):
	15	+ input_element = torch.randn(size=shape, dtype=eval('torch.' + dtype), device=torch.device("npu")) * 2000
	16	+ std_argmax = self.op_calc(input_element, dim)
	17	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	18	+ inductor_argmax = compiled_op_calc(input_element, dim)
	19	+ self.assertEqual(std_argmax, inductor_argmax, atol=1e-2, rtol=1e-2)
	20	+
	21	+instantiate_parametrized_tests(TestMaxWithIndex)
	22	+
	23	+if __name__ == "__main__":
	24	+ run_tests()

Atest/_inductor/test_arrange.py+31-0

@@ -0,0 +1,31 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestArrange(TestUtils):
	8	+ def op_calc(self, start, end, step):
	9	+ a = torch.arange(start, end, step, device=torch.device('npu'))
	10	+ y = a + a
	11	+ return y
	12	+
	13	+ @parametrize('shape', [(2, )])
	14	+ @parametrize('dtype', TestUtils._test_dtypes)
	15	+ def test_pointwise_cases(self, shape, dtype):
	16	+ s = self._generate_tensor(shape, dtype)
	17	+ start = min(s)
	18	+ end = max(s)
	19	+ step = (end - start) / 32
	20	+
	21	+ std_arrange = self.op_calc(start, end, step)
	22	+
	23	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	24	+ inductor_arrange = compiled_op_calc(start, end, step)
	25	+
	26	+ self.assertEqual(std_arrange, inductor_arrange)
	27	+
	28	+instantiate_parametrized_tests(TestArrange)
	29	+
	30	+if __name__ == "__main__":
	31	+ run_tests()

Atest/_inductor/test_attncp.py+32-0

@@ -0,0 +1,32 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestAttnCp(TestUtils):
	8	+ shape = (8, 8, 256, 128)
	9	+ dim = -1
	10	+
	11	+ def foo(self, a, b, c):
	12	+ y = a + b
	13	+ y = y.sum(self.dim)
	14	+ y = y.unsqueeze(self.dim)
	15	+ y = y.broadcast_to(self.shape) + b
	16	+ y = c + y.permute(0, 1, 3, 2)
	17	+ return y
	18	+
	19	+
	20	+ def test_pointwise_cases(self):
	21	+ a, b = [torch.randn(self.shape, dtype=torch.float32, device="npu") for _ in range(2)]
	22	+ d = torch.randn(self.shape, dtype=torch.float32, device="npu")
	23	+ c = d.permute(0, 1, 3, 2).contiguous()
	24	+ func = torch.compile(self.foo, backend="inductor")
	25	+ r = func(a, b, c)
	26	+ r1 = self.foo(a, b, c)
	27	+ self.assertEqual(r, r1, atol=1e-3, rtol=1e-3)
	28	+
	29	+instantiate_parametrized_tests(TestAttnCp)
	30	+
	31	+if __name__ == "__main__":
	32	+ run_tests()

Atest/_inductor/test_batch_norm.py+46-0

@@ -0,0 +1,46 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestNativeBatchNorm(TestUtils):
	8	+ def op_calc(self, input_element):
	9	+ # 创建权重和偏置张量
	10	+ weight = torch.ones(32).npu()
	11	+ bias = torch.zeros(32).npu()
	12	+
	13	+ # 创建运行均值和方差张量
	14	+ running_mean = torch.zeros(32).npu()
	15	+ running_var = torch.ones(32).npu()
	16	+ momentum = 0.1
	17	+ eps = 1e-05
	18	+ # 执行批量归一化
	19	+ output, running_mean_out, running_var_out = torch.native_batch_norm(
	20	+ input=input_element,
	21	+ weight=weight,
	22	+ bias=bias,
	23	+ running_mean=running_mean,
	24	+ running_var=running_var,
	25	+ training=True,
	26	+ momentum=momentum,
	27	+ eps=eps
	28	+ )
	29	+ return output, running_mean_out, running_var_out
	30	+
	31	+ @parametrize('shape', [(16, 32, 64)])
	32	+ @parametrize('dtype', ['float32'])
	33	+ def test_reduction_cases_shapes(self, shape, dtype):
	34	+ input_element = self._generate_tensor(shape, dtype)
	35	+
	36	+ std_ret, _, _ = self.op_calc(input_element)
	37	+
	38	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	39	+ inductor_ret, _, _ = compiled_op_calc(input_element)
	40	+ self.assertEqual(std_ret, inductor_ret, atol=1e-1, rtol=1e-1, equal_nan=True)
	41	+
	42	+
	43	+instantiate_parametrized_tests(TestNativeBatchNorm)
	44	+
	45	+if __name__ == "__main__":
	46	+ run_tests()

Atest/_inductor/test_broadcast.py+39-0

@@ -0,0 +1,39 @@
	1	+import copy
	2	+import torch
	3	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	4	+from testutils import TestUtils
	5	+import torch_npu
	6	+
	7	+
	8	+class TestBroadcast(TestUtils):
	9	+ broadcast_size = 128
	10	+
	11	+ def op_calc(self, a, b, dim, new_shape):
	12	+ a = a.unsqueeze(dim)
	13	+ a = a.broadcast_to(new_shape)
	14	+ b = b.unsqueeze(dim)
	15	+ b = b.broadcast_to(new_shape)
	16	+ y = a + b
	17	+ return y
	18	+
	19	+
	20	+ @parametrize('shape', [(8, 8, 256)])
	21	+ @parametrize('dtype', ['float32', 'int32', 'float16', 'bfloat16'])
	22	+ def test_view_cases(self, shape, dtype):
	23	+ a = self._generate_tensor(shape, dtype)
	24	+ b = self._generate_tensor(shape, dtype)
	25	+
	26	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	27	+ for dim in [3, 2, 1, 0]:
	28	+ new_shape = list(copy.deepcopy(shape))
	29	+ new_shape.insert(dim, self.broadcast_size)
	30	+ std_broadcast = self.op_calc(a, b, dim, new_shape)
	31	+ inductor_broadcast = compiled_op_calc(a, b, dim, new_shape)
	32	+
	33	+ self.assertEqual(std_broadcast.float(), inductor_broadcast.float(), atol=1e-3, rtol=1e-3)
	34	+
	35	+
	36	+instantiate_parametrized_tests(TestBroadcast)
	37	+
	38	+if __name__ == "__main__":
	39	+ run_tests()

Atest/_inductor/test_cat.py+27-0

@@ -0,0 +1,27 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestCat(TestUtils):
	8	+
	9	+ def op_calc(self, input_element, dim):
	10	+ return torch.cat([input_element, input_element], dim)
	11	+
	12	+ # case：change shapes
	13	+ @parametrize('shape', [(8, 16, 32, 64)])
	14	+ @parametrize('dim', [-1])
	15	+ @parametrize('dtype', ['bfloat16'])
	16	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	17	+ input_element = self._generate_tensor(shape, dtype)
	18	+ std_cat = self.op_calc(input_element, dim)
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_cat = compiled_op_calc(input_element, dim)
	21	+ self.assertEqual(std_cat, inductor_cat, atol=1e-1, rtol=1e-1, equal_nan=True)
	22	+
	23	+
	24	+instantiate_parametrized_tests(TestCat)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_ceil.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestRelu(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.ceil(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float32', 'float16', 'bfloat16'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+
	22	+ self.assertEqual(std_result, inductor_result)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestRelu)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_check_accuracy.py+70-0

@@ -0,0 +1,70 @@
	1	+import os
	2	+from unittest.mock import patch
	3	+
	4	+import torch
	5	+import torch.nn.functional as F
	6	+from torch.testing._internal.common_utils import run_tests
	7	+from testutils import TestUtils
	8	+import torch_npu
	9	+
	10	+os.environ["INDUCTOR_ASCEND_CHECK_ACCURACY"] = "1"
	11	+
	12	+
	13	+class TestCheckAccuracy(TestUtils):
	14	+ def test_check_accuracy_1(self):
	15	+ count_data_dump = 0
	16	+ count_check_accuracy = 0
	17	+
	18	+ def run(x, y):
	19	+ return F.relu(x) - y
	20	+
	21	+ from torch_npu._inductor.npu_triton_heuristics import NPUCachingAutotuner
	22	+ src_data_dump = NPUCachingAutotuner.data_dump
	23	+
	24	+ def wrap_data_dump(self, args, *kwargs):
	25	+ status = src_data_dump(self, args, *kwargs)
	26	+ if status:
	27	+ nonlocal count_data_dump
	28	+ count_data_dump += 1
	29	+ return status
	30	+
	31	+ src_check_accuracy = NPUCachingAutotuner.check_accuracy
	32	+
	33	+ def wrap_check_accuracy(self, args, *kwargs):
	34	+ status = src_check_accuracy(self, args, *kwargs)
	35	+ if status:
	36	+ nonlocal count_check_accuracy
	37	+ count_check_accuracy += 1
	38	+ return status
	39	+
	40	+ x = torch.randn(10).npu()
	41	+ y = torch.randn(10).npu()
	42	+ g = run(x, y)
	43	+
	44	+ run = torch.compile(run)
	45	+ # compile warmup
	46	+ _ = run(x, y)
	47	+
	48	+ with patch.object(NPUCachingAutotuner, "data_dump", wrap_data_dump), \
	49	+ patch.object(NPUCachingAutotuner, "check_accuracy", wrap_check_accuracy):
	50	+ self.assertTrue(torch_npu._inductor.config.dump_fx_graph)
	51	+ self.assertTrue(torch_npu._inductor.config.check_accuracy)
	52	+
	53	+ # Try run custom path and make sure no data_dump and check_accuracy is invoked.
	54	+ torch_npu._inductor.config.dump_fx_graph = False
	55	+ torch_npu._inductor.config.check_accuracy = False
	56	+ z = run(x, y)
	57	+ self.assertEqual(count_data_dump, 0)
	58	+ self.assertEqual(count_check_accuracy, 0)
	59	+ self.assertEqual(z, g)
	60	+
	61	+ torch_npu._inductor.config.dump_fx_graph = True
	62	+ torch_npu._inductor.config.check_accuracy = True
	63	+ z = run(x, y)
	64	+ self.assertEqual(count_data_dump, 1)
	65	+ self.assertEqual(count_check_accuracy, 1)
	66	+ self.assertEqual(z, g)
	67	+
	68	+
	69	+if __name__ == "__main__":
	70	+ run_tests()

Atest/_inductor/test_clamp.py+88-0

@@ -0,0 +1,88 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestClamp(TestUtils):
	8	+
	9	+ def op_calc(self, arg, min_value=None, max_value=None):
	10	+ return arg.clamp(min_value, max_value)
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases_minmax_is_tensor(self, shape, dtype):
	15	+ min_0 = self._generate_tensor(shape, dtype)
	16	+ max_0 = self._generate_tensor(shape, dtype)
	17	+
	18	+ first_element = self._generate_tensor(shape, dtype)
	19	+
	20	+ std_result = self.op_calc(first_element, min_value=min_0, max_value=max_0)
	21	+
	22	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	23	+ inductor_result = compiled_op_calc(first_element, min_value=min_0, max_value=max_0)
	24	+
	25	+ self.assertEqual(std_result, inductor_result)
	26	+
	27	+ @parametrize('shape', [(1,)])
	28	+ @parametrize('dtype', ['float32'])
	29	+ def test_pointwise_cases_single_scalar(self, shape, dtype):
	30	+ min_numel = 0
	31	+ max_numel = 100
	32	+
	33	+ first_element = 200 * torch.rand(size=shape, dtype=eval('torch.' + dtype), device=torch.device("npu"))
	34	+
	35	+ std_result = self.op_calc(first_element, min_value=min_numel, max_value=max_numel)
	36	+
	37	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	38	+ inductor_result = compiled_op_calc(first_element, min_value=min_numel, max_value=max_numel)
	39	+ self.assertEqual(std_result, inductor_result)
	40	+
	41	+ @parametrize('shape', [(1024, 32)])
	42	+ @parametrize('dtype', ['int32'])
	43	+ def test_pointwise_cases_minmax_is_number(self, shape, dtype):
	44	+ min_numel = 0
	45	+ max_numel = 100
	46	+
	47	+ first_element = self._generate_tensor(shape, dtype)
	48	+
	49	+ std_result = self.op_calc(first_element, min_value=min_numel, max_value=max_numel)
	50	+
	51	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	52	+ inductor_result = compiled_op_calc(first_element, min_value=min_numel, max_value=max_numel)
	53	+
	54	+ self.assertEqual(std_result, inductor_result)
	55	+
	56	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	57	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	58	+ def test_pointwise_cases_max_only(self, shape, dtype):
	59	+ max_numel = 100
	60	+
	61	+ first_element = self._generate_tensor(shape, dtype)
	62	+
	63	+ std_result = self.op_calc(first_element, min_value=None, max_value=max_numel)
	64	+
	65	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	66	+ inductor_result = compiled_op_calc(first_element, min_value=None, max_value=max_numel)
	67	+
	68	+ self.assertEqual(std_result, inductor_result)
	69	+
	70	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	71	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	72	+ def test_pointwise_cases_min_only(self, shape, dtype):
	73	+ min_numel = 0
	74	+
	75	+ first_element = self._generate_tensor(shape, dtype)
	76	+
	77	+ std_result = self.op_calc(first_element, min_value=min_numel, max_value=None)
	78	+
	79	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	80	+ inductor_result = compiled_op_calc(first_element, min_value=min_numel, max_value=None)
	81	+
	82	+ self.assertEqual(std_result, inductor_result)
	83	+
	84	+
	85	+instantiate_parametrized_tests(TestClamp)
	86	+
	87	+if __name__ == "__main__":
	88	+ run_tests()

Atest/_inductor/test_clone.py+27-0

@@ -0,0 +1,27 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestClone(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return torch.clone(input_element)
	10	+
	11	+ @parametrize('shape', [(8, 64, 128)])
	12	+ @parametrize('dim', [0])
	13	+ @parametrize('dtype', ['float32'])
	14	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	15	+ input_element = self._generate_tensor(shape, dtype)
	16	+ std_ret = self.op_calc(input_element, dim)
	17	+
	18	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	19	+ inductor_ret = compiled_op_calc(input_element, dim)
	20	+
	21	+ self.assertEqual(std_ret, inductor_ret, equal_nan=True)
	22	+
	23	+
	24	+instantiate_parametrized_tests(TestClone)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_codecache.py+27-0

@@ -0,0 +1,27 @@
	1	+import pytest
	2	+import torch
	3	+from torch.testing._internal.common_utils import run_tests
	4	+from torch._inductor.codecache import CacheBase
	5	+from testutils import TestUtils
	6	+import torch_npu
	7	+import torch_npu._inductor
	8	+
	9	+
	10	+class TestCodeCache(TestUtils):
	11	+ def test_codecache(self):
	12	+ device_properties = torch_npu.npu.get_device_properties(
	13	+ torch_npu.npu.current_device()
	14	+ )
	15	+
	16	+ system1 = CacheBase.get_system()
	17	+ self.assertEqual(system1["device"]["name"], device_properties.name)
	18	+ self.assertEqual(system1["version"]["cann"], torch.version.cann)
	19	+
	20	+ from torch_npu.contrib import transfer_to_npu
	21	+ system2 = CacheBase.get_system()
	22	+ self.assertEqual(system2["device"]["name"], device_properties.name)
	23	+ self.assertEqual(system2["version"]["cann"], torch.version.cann)
	24	+
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_cos.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestLog(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.cos(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+
	22	+ self.assertEqual(std_result, inductor_result)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestLog)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_debug_msg.py+129-0

@@ -0,0 +1,129 @@
	1	+
	2	+import os
	3	+import re
	4	+import logging
	5	+import tempfile
	6	+from pathlib import Path
	7	+import torch
	8	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	9	+from torch._inductor import config
	10	+from testutils import TestUtils
	11	+import torch_npu
	12	+
	13	+os.environ["INDUCTOR_ASCEND_DUMP_FX_GRAPH"] = "1"
	14	+os.environ["TORCH_COMPILE_DEBUG"] = "1"
	15	+
	16	+
	17	+class TestDebugMsg(TestUtils):
	18	+ @parametrize('shape_x', [(32, 512, 64)])
	19	+ @parametrize('shape_y', [(32, 1, 64)])
	20	+ @parametrize('dtype', ['float32'])
	21	+ def test_case1(self, shape_x, shape_y, dtype):
	22	+ x = self._generate_tensor(shape_x, dtype)
	23	+ y = self._generate_tensor(shape_y, dtype)
	24	+
	25	+
	26	+ def run_case1(x, y):
	27	+ z = x + y
	28	+ return z
	29	+
	30	+ run = torch.compile(run_case1, backend='inductor')
	31	+ with config.patch(
	32	+ {
	33	+ "trace.debug_dir": tempfile.mkdtemp(),
	34	+ "force_disable_caches": True,
	35	+ }
	36	+ ):
	37	+ with self.assertLogs(
	38	+ logging.getLogger("torch._inductor.debug"), level=logging.WARNING
	39	+ ) as cm:
	40	+ run(x, y)
	41	+
	42	+ self.assertEqual(len(cm.output), 1)
	43	+ m = re.match(r"WARNING.* debug trace: (.*)", cm.output[0])
	44	+ self.assertTrue(m)
	45	+ filename = Path(m.group(1))
	46	+ self.assertTrue(filename.is_dir())
	47	+ content = open(filename / "output_code.py").read().rstrip()
	48	+
	49	+ self.assertIn(
	50	+ "# SchedulerNodes: [SchedulerNode(name='op0')]",
	51	+ content
	52	+ )
	53	+
	54	+ self.assertIn(
	55	+ """
	56	+# def forward(self, arg0_1, arg1_1):
	57	+# expand = torch.ops.aten.expand.default(arg1_1, [32, 512, 64]); arg1_1 = None
	58	+# add = torch.ops.aten.add.Tensor(arg0_1, expand); arg0_1 = expand = None
	59	+# return (add,)""",
	60	+ content
	61	+ )
	62	+
	63	+ self.assertIn(
	64	+ """
	65	+# inputs: [FakeTensor(..., device='npu:0', size=(32, 512, 64), strides=(32768, 64, 1)), FakeTensor(..., device='npu:0', size=(32, 1, 64), strides=(64, 64, 1))]
	66	+# outputs: [FakeTensor(..., device='npu:0', size=(32, 512, 64), strides=(32768, 64, 1))]""",
	67	+ content
	68	+ )
	69	+
	70	+
	71	+ @parametrize('shape_x', [(32, 512, 64)])
	72	+ @parametrize('shape_y', [(32, 1, 64)])
	73	+ @parametrize('dtype', ['float32'])
	74	+ def test_case2(self, shape_x, shape_y, dtype):
	75	+ x = self._generate_tensor(shape_x, dtype)
	76	+ y = self._generate_tensor(shape_y, dtype)
	77	+
	78	+
	79	+ def run_case2(x, y):
	80	+ z = x + y
	81	+ z = z.repeat([256, 1, 1])
	82	+ return z
	83	+
	84	+ run = torch.compile(run_case2, backend='inductor')
	85	+ with config.patch(
	86	+ {
	87	+ "trace.debug_dir": tempfile.mkdtemp(),
	88	+ "force_disable_caches": True,
	89	+ }
	90	+ ):
	91	+ with self.assertLogs(
	92	+ logging.getLogger("torch._inductor.debug"), level=logging.WARNING
	93	+ ) as cm:
	94	+ run(x, y)
	95	+
	96	+ self.assertEqual(len(cm.output), 1)
	97	+ m = re.match(r"WARNING.* debug trace: (.*)", cm.output[0])
	98	+ self.assertTrue(m)
	99	+ filename = Path(m.group(1))
	100	+ self.assertTrue(filename.is_dir())
	101	+ content = open(filename / "output_code.py").read().rstrip()
	102	+
	103	+ self.assertIn(
	104	+ "# SchedulerNodes: [SchedulerNode(name='op0')]",
	105	+ content
	106	+ )
	107	+
	108	+ self.assertIn(
	109	+ """
	110	+# def forward(self, arg0_1, arg1_1):
	111	+# expand = torch.ops.aten.expand.default(arg1_1, [32, 512, 64]); arg1_1 = None
	112	+# add = torch.ops.aten.add.Tensor(arg0_1, expand); arg0_1 = expand = None
	113	+# repeat = torch.ops.aten.repeat.default(add, [256, 1, 1]); add = None
	114	+# return (repeat,)""",
	115	+ content
	116	+ )
	117	+
	118	+ self.assertIn(
	119	+ """
	120	+# inputs: [FakeTensor(..., device='npu:0', size=(32, 512, 64), strides=(32768, 64, 1)), FakeTensor(..., device='npu:0', size=(32, 1, 64), strides=(64, 64, 1))]
	121	+# outputs: [FakeTensor(..., device='npu:0', size=(8192, 512, 64), strides=(32768, 64, 1))]""",
	122	+ content
	123	+ )
	124	+
	125	+
	126	+instantiate_parametrized_tests(TestDebugMsg)
	127	+
	128	+if __name__ == "__main__":
	129	+ run_tests()

Atest/_inductor/test_device_put.py+34-0

@@ -0,0 +1,34 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestDevicePut(TestUtils):
	8	+ def op_calc(self, input_element1, input_element2):
	9	+ return torch.add(input_element1, input_element2)
	10	+
	11	+ @parametrize('shape', [(8, 16, 8)])
	12	+ @parametrize('dtype', ['int32'])
	13	+ def test_cases_shapes(self, shape, dtype):
	14	+ low = 0
	15	+ high = 2
	16	+ dtype = eval('torch.' + dtype)
	17	+ npu_device = torch.device('npu:0')
	18	+ input_element1_tmp = torch.randint(low, high, shape, dtype=dtype).cpu()
	19	+ input_element2_tmp = torch.randint(low, high, shape, dtype=dtype).cpu()
	20	+ input_element1 = torch.ops.prims.device_put(input_element1_tmp, npu_device)
	21	+ input_element2 = torch.ops.prims.device_put(input_element2_tmp, npu_device)
	22	+
	23	+ std_ret = self.op_calc(input_element1, input_element2)
	24	+
	25	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	26	+ inductor_ret = compiled_op_calc(input_element1, input_element2)
	27	+
	28	+ self.assertEqual(std_ret, inductor_ret)
	29	+
	30	+
	31	+instantiate_parametrized_tests(TestDevicePut)
	32	+
	33	+if __name__ == "__main__":
	34	+ run_tests()

Atest/_inductor/test_div.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestDiv(TestUtils):
	8	+ def op_calc(self, first_element, second_element):
	9	+ result = torch.div(first_element, second_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+ second_element = self._generate_tensor(shape, dtype)
	17	+
	18	+ std_result = self.op_calc(first_element, second_element)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_result = compiled_op_calc(first_element, second_element)
	22	+ self.assertEqual(std_result, inductor_result)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestDiv)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_embedding.py+41-0

@@ -0,0 +1,41 @@
	1	+import torch
	2	+import torch.nn as nn
	3	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	4	+from testutils import TestUtils
	5	+import torch_npu
	6	+
	7	+
	8	+class TestEmbeddingDense(TestUtils):
	9	+ def op_calc(self, arg, embedding):
	10	+ output = embedding(arg)
	11	+ return output
	12	+
	13	+ # UT skip, reason: precision fail
	14	+ # Added to pytorch-disable-tests.json
	15	+ def test_pointwise_cases(self):
	16	+
	17	+ arg0 = torch.tensor([[14, 1, 2, 10, 0, 10, 0],
	18	+ [9, 13, 13, 4, 7, 15, 14],
	19	+ [8, 0, 3, 15, 4, 2, 6],
	20	+ [15, 12, 13, 9, 0, 8, 1],
	21	+ [8, 15, 4, 15, 12, 9, 3],
	22	+ [6, 11, 12, 8, 0, 13, 8],
	23	+ [4, 10, 1, 12, 0, 0, 4],
	24	+ [6, 6, 15, 6, 0, 10, 15],
	25	+ [2, 5, 14, 0, 5, 7, 9],
	26	+ [13, 4, 14, 11, 11, 9, 2],
	27	+ [1, 1, 5, 1, 1, 6, 14],
	28	+ [3, 9, 8, 4, 13, 8, 3],
	29	+ [4, 10, 8, 13, 6, 8, 3]], device='npu:0')
	30	+ embedding = nn.Embedding(16, 128).npu()
	31	+ std_sub = self.op_calc(arg0, embedding)
	32	+
	33	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	34	+ inductor_sum = compiled_op_calc(arg0, embedding)
	35	+ self.assertEqual(std_sub, inductor_sum)
	36	+
	37	+
	38	+instantiate_parametrized_tests(TestEmbeddingDense)
	39	+
	40	+if __name__ == "__main__":
	41	+ run_tests()

Atest/_inductor/test_embedding_fallback.py+29-0

@@ -0,0 +1,29 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestEmbeddingDenseBackward(TestUtils):
	8	+ def op_calc(self, slice_4, sum_23):
	9	+ result = torch.ops.aten.embedding_dense_backward.default(sum_23, slice_4, 512, -1, False)
	10	+ return result
	11	+
	12	+ @parametrize('shape', [(1, 512, 128)])
	13	+ @parametrize('dtype', ['float32'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = torch.randint(low=0, high=128, size=(1, 512), dtype=torch.int64).npu()
	16	+ second_element = self._generate_tensor(shape, dtype)
	17	+
	18	+ std_result = self.op_calc(first_element, second_element)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_result = compiled_op_calc(first_element, second_element)
	22	+
	23	+ self.assertEqual(std_result, inductor_result, atol=1e-1, rtol=1e-1)
	24	+
	25	+
	26	+instantiate_parametrized_tests(TestEmbeddingDenseBackward)
	27	+
	28	+if __name__ == "__main__":
	29	+ run_tests()

Atest/_inductor/test_empty.py+45-0

@@ -0,0 +1,45 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestEmpty(TestUtils):
	8	+ def op_calc(self):
	9	+ x = torch.empty(8, 64, 128, dtype=torch.float32).npu()
	10	+ x.uniform_(-100, 100)
	11	+ return x
	12	+
	13	+ def op_calc_empty_permuted(self):
	14	+ input_shape = (8, 64, 128)
	15	+ physical_layout = (0, 1, 2)
	16	+ x = torch.empty_permuted(input_shape, physical_layout).npu()
	17	+ x.uniform_(-100, 100)
	18	+ return x
	19	+
	20	+ # case： change shapes
	21	+ @parametrize('shape', [(8, 64, 128)])
	22	+ @parametrize('dim', [0])
	23	+ @parametrize('dtype', ['float32'])
	24	+ def test_cases_empty(self, shape, dim, dtype):
	25	+
	26	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	27	+ inductor_ret = compiled_op_calc()
	28	+
	29	+ self.assertTrue(inductor_ret.numel() > 0)
	30	+
	31	+ @parametrize('shape', [(8, 64, 128)])
	32	+ @parametrize('dim', [0])
	33	+ @parametrize('dtype', ['float32'])
	34	+ def test_cases_empty_permuted(self, shape, dim, dtype):
	35	+ compiled_op_calc = torch.compile(self.op_calc_empty_permuted, backend="inductor")
	36	+ inductor_ret = compiled_op_calc()
	37	+
	38	+ self.assertTrue(inductor_ret.numel() > 0)
	39	+
	40	+
	41	+instantiate_parametrized_tests(TestEmpty)
	42	+
	43	+if __name__ == "__main__":
	44	+ run_tests()
	45	+

Atest/_inductor/test_eq.py+35-0

@@ -0,0 +1,35 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestEq(TestUtils):
	8	+ def op_calc(self, first_element, second_element):
	9	+ return torch.eq(first_element, second_element)
	10	+
	11	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	12	+ @parametrize('dtype', ['float32', 'int32', 'float16', 'bfloat16'])
	13	+ def test_pointwise_cases(self, shape, dtype):
	14	+
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+ second_element = first_element.clone()
	17	+
	18	+ # randomly change some elements in second tensor
	19	+ flat_second_view = second_element.flatten()
	20	+ num_elements_to_change = first_element.numel() // 3
	21	+ random_indices = torch.randint(0, first_element.numel(), (num_elements_to_change,))
	22	+ flat_second_view[random_indices] = 1 - flat_second_view[random_indices]
	23	+
	24	+ std_result = self.op_calc(first_element, second_element)
	25	+
	26	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	27	+ inductor_result = compiled_op_calc(first_element, second_element)
	28	+
	29	+ self.assertEqual(std_result, inductor_result)
	30	+
	31	+
	32	+instantiate_parametrized_tests(TestEq)
	33	+
	34	+if __name__ == "__main__":
	35	+ run_tests()

Atest/_inductor/test_exceptions.py+71-0

@@ -0,0 +1,71 @@
	1	+import functools
	2	+from functools import partial
	3	+
	4	+from testutils import TestUtils
	5	+import torch
	6	+from torch._inductor.codecache import _load_triton_kernel_from_source
	7	+from torch.testing._internal.common_utils import run_tests
	8	+import torch_npu
	9	+
	10	+
	11	+src_code_1 = '''
	12	+import triton
	13	+import triton.language as tl
	14	+from triton.compiler.compiler import AttrsDescriptor
	15	+
	16	+from torch._inductor.runtime import triton_helpers, triton_heuristics
	17	+from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
	18	+from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, DeviceProperties
	19	+
	20	+from torch._inductor.runtime import triton_helpers
	21	+from torch_npu._inductor import npu_triton_heuristics
	22	+from torch_npu._inductor import npu_triton_helpers
	23	+from torch_npu._inductor.runtime import NPUDeviceProperties
	24	+from torch_npu._inductor.npu_triton_helpers import libdevice, math as tl_math
	25	+import torch
	26	+import torch_npu
	27	+
	28	+@npu_triton_heuristics.pointwise_npu_index(
	29	+ size_hints=[16384, 32], tile_hint=TileHint.DEFAULT,
	30	+ filename=__file__,
	31	+ triton_meta={'signature': {'in_ptr0': 'fp16', 'in_ptr1': 'fp16', 'out_ptr0': '*fp16', 'y0_numel': 'i32', 'x1_numel': 'i32'},
	32	+ 'device': NPUDeviceProperties(type='npu', index=0, multi_processor_count=40, cc='Ascend910B3',
	33	+ major=None, regs_per_multiprocessor=None, max_threads_per_multi_processor=None, warp_size=32),
	34	+ 'constants': {}, 'mix_mode': 'aiv'},
	35	+ inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_unk_fused_add_0', 'mutated_arg_names': [],
	36	+ 'backend_hash': 'bc71dba4086164e7ac2b0779fa861dbf7467f0265d4a57b8f48cf6dda02b150f', 'split_axis': [0],
	37	+ 'tiling_axis': [0, 1], 'axis_names': ['y0', 'x1'], 'low_dims': {1}, 'numof_reduction_axis': 0,
	38	+ 'split_axis_dtype': torch.float16, 'dual_reduction': False, 'traced_graph_hash': 'TRACED_GRAPH_HASH',
	39	+ 'traced_graph_dir': 'TRACED_GRAPH_DIR'},
	40	+ min_elem_per_thread=0
	41	+)
	42	+@triton.jit
	43	+def triton_unk_fused_add_0(in_ptr0, in_ptr1, out_ptr0, y0_numel, x1_numel, Y0BLOCK: tl.constexpr, Y0BLOCK_SUB: tl.constexpr, X1BLOCK_SUB: tl.constexpr):
	44	+ y0_offset = tl.program_id(0) * Y0BLOCK
	45	+ base_y0= tl.arange(0, Y0BLOCK_SUB)
	46	+ loops_y0 = (Y0BLOCK + Y0BLOCK_SUB - 1) // Y0BLOCK_SUB
	47	+ base_x1= tl.arange(0, X1BLOCK_SUB)
	48	+ loops_x1 = (x1_numel + X1BLOCK_SUB - 1) // X1BLOCK_SUB
	49	+ for loop_y0 in range(loops_y0):
	50	+ y0 = y0_offset + (loop_y0 * Y0BLOCK_SUB) + base_y0[:,None]
	51	+ y0_mask = y0 < min(Y0BLOCK+y0_offset, y0_numel)
	52	+ for loop_x1 in range(loops_x1):
	53	+ x1 = (loop_x1 * X1BLOCK_SUB) + base_x1[None,:]
	54	+ x1_mask = x1 < x1_numel
	55	+ tmp0 = tl.load(in_ptr0 + (x1 + 128*y0), x1_mask & y0_mask)
	56	+ # Not define tmp1 and make error manually for triton: 'tmp1 is not defined'
	57	+ tmp2 = tmp0 + tmp1
	58	+ tl.store(out_ptr0 + (x1 + 32*y0), tmp2, x1_mask & y0_mask)
	59	+'''
	60	+
	61	+
	62	+class TestExceptions(TestUtils):
	63	+ def test_triton_kernel_failed(self):
	64	+ with self.assertRaisesRegex(Exception, "tmp1 is not defined"):
	65	+ load_kernel = functools.partial(_load_triton_kernel_from_source, "triton_unk_fused_add_0", src_code_1)
	66	+ kernel = load_kernel()
	67	+ kernel.precompile()
	68	+
	69	+
	70	+if __name__ == "__main__":
	71	+ run_tests()

Atest/_inductor/test_exp.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestExp(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.exp(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+ self.assertEqual(std_result, inductor_result, atol=1e-1, rtol=1e-1, equal_nan=True)
	22	+
	23	+
	24	+instantiate_parametrized_tests(TestExp)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()
	28	+

Atest/_inductor/test_expm1.py+27-0

@@ -0,0 +1,27 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestExpm1(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.expm1(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	19	+ inductor_result = compiled_op_calc(first_element)
	20	+
	21	+ self.assertEqual(std_result, inductor_result, atol=1e-3, rtol=1e-3, equal_nan=True)
	22	+
	23	+
	24	+instantiate_parametrized_tests(TestExpm1)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_floor.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestFloor(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.floor(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+
	22	+ self.assertEqual(std_result, inductor_result)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestFloor)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_force_fallback.py+57-0

@@ -0,0 +1,57 @@
	1	+import os
	2	+from unittest.mock import patch
	3	+
	4	+import torch
	5	+import torch.nn.functional as F
	6	+from torch.testing._internal.common_utils import run_tests
	7	+from testutils import TestUtils
	8	+import torch_npu
	9	+
	10	+os.environ["INDUCTOR_ASCEND_DUMP_FX_GRAPH"] = "1"
	11	+
	12	+
	13	+class TestForceFallback(TestUtils):
	14	+ def test_case1(self):
	15	+ op_list = []
	16	+
	17	+ def opoverload_call(self, /, args, *kwargs):
	18	+ op_list.append(str(self))
	19	+ return self._op(args, *kwargs)
	20	+
	21	+ def run(x, y):
	22	+ return F.relu(x) + y
	23	+
	24	+ x = torch.randn(10).npu()
	25	+ y = torch.randn(10).npu()
	26	+ g = run(x, y)
	27	+
	28	+ run = torch.compile(run)
	29	+ # compile warmup
	30	+ _ = run(x, y)
	31	+
	32	+ with patch.object(torch._ops.OpOverload, "__call__", opoverload_call):
	33	+ op_list.clear()
	34	+ z = run(x, y)
	35	+ self.assertTrue(len(op_list) == 0)
	36	+ self.assertEqual(z, g)
	37	+
	38	+ op_list.clear()
	39	+ torch_npu._inductor.config.force_fallback_kernel_id = [0]
	40	+ z = run(x, y)
	41	+ self.assertTrue("aten.relu.default" in op_list)
	42	+ self.assertTrue("aten.add.Tensor" in op_list)
	43	+ self.assertEqual(z, g)
	44	+
	45	+ op_list.clear()
	46	+ torch_npu._inductor.config.force_fallback_kernel_id = 'all'
	47	+ z = run(x, y)
	48	+ self.assertTrue("aten.relu.default" in op_list)
	49	+ self.assertTrue("aten.add.Tensor" in op_list)
	50	+ self.assertEqual(z, g)
	51	+
	52	+ # reset
	53	+ torch_npu._inductor.config.force_fallback_kernel_id = []
	54	+
	55	+
	56	+if __name__ == "__main__":
	57	+ run_tests()

Atest/_inductor/test_ge.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestGe(TestUtils):
	8	+ def op_calc(self, first_element, second_element):
	9	+ return torch.ge(first_element, second_element)
	10	+
	11	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	12	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32'])
	13	+ def test_pointwise_cases(self, shape, dtype):
	14	+ first_element = self._generate_tensor(shape, dtype)
	15	+ second_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element, second_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element, second_element)
	21	+
	22	+ self.assertEqual(std_result, inductor_result)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestGe)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_geometric.py+36-0

@@ -0,0 +1,36 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestGeometric(TestUtils):
	8	+ def op_calc(self):
	9	+ # 创建一个形状为 (3, 3)的张量，每个位置的概率为 0.5
	10	+ prob = torch.full((16, 16), 0.5).npu()
	11	+
	12	+ #使用 aten.geometric生成几何分布的随机数
	13	+ geometric_tensor = torch.ops.aten.geometric(prob, p=0.5)
	14	+
	15	+ return geometric_tensor
	16	+
	17	+ # UT skip, reason: this has problem in torch 260
	18	+ # Added to pytorch-disable-tests.json
	19	+ @parametrize('shape', [(16, 16, 16)])
	20	+ @parametrize('dim', [0])
	21	+ @parametrize('dtype', ['int32'])
	22	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	23	+ std_ret = self.op_calc()
	24	+ std_ret_mean = torch.mean(std_ret)
	25	+
	26	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	27	+ inductor_ret = compiled_op_calc()
	28	+
	29	+ inductor_ret_mean = torch.mean(inductor_ret)
	30	+ self.assertTrue(inductor_ret_mean is not None)
	31	+
	32	+
	33	+instantiate_parametrized_tests(TestGeometric)
	34	+
	35	+if __name__ == "__main__":
	36	+ run_tests()

Atest/_inductor/test_gt.py+30-0

@@ -0,0 +1,30 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestGt(TestUtils):
	8	+ def op_calc(self, first_element, second_element):
	9	+ result = torch.gt(first_element, second_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+ second_element = self._generate_tensor(shape, dtype)
	17	+
	18	+ std_result = self.op_calc(first_element, second_element)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_result = compiled_op_calc(first_element, second_element)
	22	+
	23	+ self.assertEqual(std_result, inductor_result)
	24	+
	25	+
	26	+instantiate_parametrized_tests(TestGt)
	27	+
	28	+if __name__ == "__main__":
	29	+ run_tests()
	30	+

Atest/_inductor/test_high_order_sum.py+25-0

@@ -0,0 +1,25 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSum(TestUtils):
	8	+ def op_sum(self, npu_dropout_backward_9):
	9	+ view_337: "f32[32768, 256]" = torch.ops.aten.view.default(npu_dropout_backward_9, [32768, 256])
	10	+ sum_63: "f32[1, 256]" = torch.ops.aten.sum.dim_IntList(view_337, [0], True)
	11	+ view_338: "f32[256]" = torch.ops.aten.view.default(sum_63, [256])
	12	+ return view_338
	13	+
	14	+
	15	+ def test_high_order_sum(self):
	16	+ npu_dropout_backward_9 = torch.randn((32768, 256), device='npu', dtype=torch.float32)
	17	+ ref = self.op_sum(npu_dropout_backward_9)
	18	+ func = torch.compile(self.op_sum, backend="inductor", dynamic=False)
	19	+ calc = func(npu_dropout_backward_9)
	20	+
	21	+ self.assertEqual(ref, calc, atol=1e-3, rtol=1e-3)
	22	+
	23	+
	24	+if __name__ == "__main__":
	25	+ run_tests()

Atest/_inductor/test_issue54.py+60-0

@@ -0,0 +1,60 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class Test_issue54(TestUtils):
	8	+ def func_layernorm(self, args):
	9	+ add_3, primals_6, primals_7, view, primals_9, permute_1, primals_10, primals_11 = args
	10	+ permute: "f32[256, 256]" = torch.ops.aten.permute.default(primals_6, [1, 0])
	11	+ addmm: "f32[32768, 256]" = torch.ops.aten.addmm.default(primals_7, view, permute)
	12	+ view_1: "f32[64, 512, 256]" = torch.ops.aten.view.default(addmm, [64, 512, 256])
	13	+ addmm_1: "f32[32768, 256]" = torch.ops.aten.addmm.default(primals_9, view, permute_1)
	14	+ view_3: "f32[64, 512, 256]" = torch.ops.aten.view.default(addmm_1, [64, 512, 256])
	15	+ view_4: "f32[64, 512, 4, 64]" = torch.ops.aten.view.default(view_3, [64, 512, 4, 64])
	16	+ permute_2: "f32[64, 4, 512, 64]" = torch.ops.aten.permute.default(view_4, [0, 2, 1, 3])
	17	+ permute_3: "f32[256, 256]" = torch.ops.aten.permute.default(primals_10, [1, 0])
	18	+ addmm_2: "f32[32768, 256]" = torch.ops.aten.addmm.default(primals_11, view, permute_3)
	19	+ view_6: "f32[64, 512, 256]" = torch.ops.aten.view.default(addmm_2, [64, 512, 256])
	20	+
	21	+ view_8: "f32[64, 512, 4, 64]" = torch.ops.aten.view.default(view_1, [64, 512, 4, 64])
	22	+ permute_5: "f32[64, 4, 512, 64]" = torch.ops.aten.permute.default(view_8, [0, 2, 1, 3])
	23	+
	24	+ permute_6: "f32[64, 4, 64, 512]" = torch.ops.aten.permute.default(permute_2, [0, 1, 3, 2])
	25	+ expand_1: "f32[64, 4, 512, 64]" = torch.ops.aten.expand.default(permute_5, [64, 4, 512, 64])
	26	+ clone: "f32[64, 4, 512, 64]" = torch.ops.aten.clone.default(expand_1, memory_format=torch.contiguous_format)
	27	+ view_9: "f32[256, 512, 64]" = torch.ops.aten.view.default(clone, [256, 512, 64])
	28	+ expand_2: "f32[64, 4, 64, 512]" = torch.ops.aten.expand.default(permute_6, [64, 4, 64, 512])
	29	+ clone_1: "f32[64, 4, 64, 512]" = torch.ops.aten.clone.default(expand_2, memory_format=torch.contiguous_format)
	30	+ view_10: "f32[256, 64, 512]" = torch.ops.aten.view.default(clone_1, [256, 64, 512])
	31	+ bmm: "f32[256, 512, 512]" = torch.ops.aten.bmm.default(view_9, view_10)
	32	+ view_7: "f32[64, 512, 4, 64]" = torch.ops.aten.view.default(view_6, [64, 512, 4, 64])
	33	+ permute_4: "f32[64, 4, 512, 64]" = torch.ops.aten.permute.default(view_7, [0, 2, 1, 3])
	34	+ expand_4: "f32[64, 4, 512, 64]" = torch.ops.aten.expand.default(permute_4, [64, 4, 512, 64])
	35	+ clone_2: "f32[64, 4, 512, 64]" = torch.ops.aten.clone.default(expand_4, memory_format=torch.contiguous_format)
	36	+ view_13: "f32[256, 512, 64]" = torch.ops.aten.view.default(clone_2, [256, 512, 64])
	37	+
	38	+ return bmm, view_13
	39	+
	40	+ def test_issue54(self):
	41	+ device = 'npu'
	42	+ add_3 = torch.randn((64, 512, 256), device=device, dtype=torch.float32)
	43	+ primals_6 = torch.randn((256, 256), device=device, dtype=torch.float32)
	44	+ primals_7 = torch.randn((256), device=device, dtype=torch.float32)
	45	+ view = torch.randn((32768, 256), device=device, dtype=torch.float32)
	46	+ primals_9 = torch.randn((256), device=device, dtype=torch.float32)
	47	+ permute_1 = torch.randn((256, 256), device=device, dtype=torch.float32)
	48	+ primals_10 = torch.randn((256, 256), device=device, dtype=torch.float32)
	49	+ primals_11 = torch.randn((256), device=device, dtype=torch.float32)
	50	+ args = (add_3, primals_6, primals_7, view, primals_9, permute_1, primals_10, primals_11)
	51	+ ref = self.func_layernorm(args)
	52	+ func = torch.compile(self.func_layernorm, backend="inductor", dynamic=False,
	53	+ options={"unroll_reductions_threshold": 1, "aggressive_fusion": True})
	54	+ calc = func(args)
	55	+ self.assertEqual(ref[0], calc[0], atol=1e-2, rtol=1e-2)
	56	+ self.assertEqual(ref[1], calc[1], atol=1e-2, rtol=1e-2)
	57	+
	58	+
	59	+if __name__ == "__main__":
	60	+ run_tests()

Atest/_inductor/test_issue57.py+35-0

@@ -0,0 +1,35 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class Test_issue57(TestUtils):
	8	+ def op_sum(self, view_12, embedding_1, slice_11):
	9	+ permute_7 = torch.ops.aten.permute.default(embedding_1, [2, 0, 1])
	10	+ embedding_1 = None
	11	+ unsqueeze_4 = torch.ops.aten.unsqueeze.default(permute_7, 0)
	12	+ permute_7 = None
	13	+
	14	+ add_5 = torch.ops.aten.add.Tensor(unsqueeze_4, slice_11)
	15	+ slice_8 = slice_11 = None
	16	+ add_6 = torch.ops.aten.add.Tensor(view_12, add_5)
	17	+ view_12 = None
	18	+ return add_6
	19	+
	20	+ def test_issue57(self):
	21	+ device = 'npu'
	22	+ embedding_1 = torch.randn((512, 512, 64), device=device, dtype=torch.float32)
	23	+ primals_221 = torch.randn((1, 1, 1, 512), device=device, dtype=torch.float32)
	24	+ view_12 = torch.randn((1, 64, 512, 512), device=device, dtype=torch.float32)
	25	+ slice_11 = torch.randn((1, 1, 1, 512), device=device, dtype=torch.float32)
	26	+
	27	+ ref = self.op_sum(view_12, embedding_1, primals_221)
	28	+ func = torch.compile(self.op_sum, backend="inductor", dynamic=False)
	29	+ calc = func(view_12, embedding_1, primals_221)
	30	+
	31	+ self.assertEqual(ref, calc, atol=1e-3, rtol=1e-3)
	32	+
	33	+
	34	+if __name__ == "__main__":
	35	+ run_tests()

Atest/_inductor/test_issue59.py+38-0

@@ -0,0 +1,38 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class Test_issue59(TestUtils):
	8	+ def layernorm_backward(self, x, y, z):
	9	+ sum_0 = torch.sum(x)
	10	+ mean = sum_0 / torch.numel(sum_0)
	11	+ sub = x - mean
	12	+ sqr = sub * sub
	13	+ sum_1 = torch.sum(sqr)
	14	+ mean_1 = sum_1 / torch.numel(sum_1) + 1e-05
	15	+ rsqrt = torch.rsqrt(mean_1)
	16	+ mul = sub * rsqrt
	17	+ mul_1 = mul * y
	18	+ add = mul_1 + z
	19	+ mean_2 = rsqrt / torch.numel(rsqrt)
	20	+ return mul, add, mean_2
	21	+
	22	+ def test_issue59(self):
	23	+ device = 'npu'
	24	+ x = torch.randn((1, 1024), device=device, dtype=torch.float32)
	25	+ y = torch.randn((1, 1024), device=device, dtype=torch.float32)
	26	+ z = torch.randn((1, 1024), device=device, dtype=torch.float32)
	27	+
	28	+ mul, add, mean_2 = self.layernorm_backward(x, y, z)
	29	+ func = torch.compile(self.layernorm_backward, backend="inductor", dynamic=False)
	30	+ mul_t, add_t, mean_2_t = func(x, y, z)
	31	+
	32	+ self.assertEqual(mul, mul_t, atol=1e-3, rtol=1e-3)
	33	+ self.assertEqual(add, add_t, atol=1e-3, rtol=1e-3)
	34	+ self.assertEqual(mean_2, mean_2_t, atol=1e-3, rtol=1e-3)
	35	+
	36	+
	37	+if __name__ == "__main__":
	38	+ run_tests()

Atest/_inductor/test_issue62.py+47-0

@@ -0,0 +1,47 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class Test_issue62(TestUtils):
	8	+ def op_func(self, addmm_5, add):
	9	+ split = torch.ops.aten.split.Tensor(addmm_5, 1536, 1)
	10	+ getitem = split[0]
	11	+ getitem_1 = split[1]
	12	+ getitem_2 = split[2]
	13	+ getitem_3 = split[3]
	14	+ getitem_4 = split[4]
	15	+ getitem_5 = split[5]
	16	+
	17	+ clone_1 = torch.ops.aten.clone.default(add, memory_format=torch.contiguous_format)
	18	+ convert_element_type_25 = torch.ops.prims.convert_element_type.default(clone_1, torch.float32)
	19	+ var_mean = torch.ops.aten.var_mean.correction(convert_element_type_25, [2], correction=0, keepdim=True)
	20	+ getitem_6 = var_mean[0]
	21	+ getitem_7 = var_mean[1]
	22	+ add_3 = torch.ops.aten.add.Tensor(getitem_6, 1e-06)
	23	+ rsqrt = torch.ops.aten.rsqrt.default(add_3)
	24	+ sub = torch.ops.aten.sub.Tensor(clone_1, getitem_7)
	25	+ mul_7 = torch.ops.aten.mul.Tensor(sub, rsqrt)
	26	+ convert_element_type_26 = torch.ops.prims.convert_element_type.default(mul_7, torch.float16)
	27	+ slice_11 = torch.ops.aten.slice.Tensor(getitem_1, 0, 0, 9223372036854775807)
	28	+ unsqueeze_2 = torch.ops.aten.unsqueeze.default(slice_11, 1)
	29	+ add_4 = torch.ops.aten.add.Tensor(unsqueeze_2, 1)
	30	+ mul_8 = torch.ops.aten.mul.Tensor(convert_element_type_26, add_4)
	31	+ slice_12 = torch.ops.aten.slice.Tensor(getitem, 0, 0, 9223372036854775807)
	32	+ unsqueeze_3 = torch.ops.aten.unsqueeze.default(slice_12, 1)
	33	+ add_5 = torch.ops.aten.add.Tensor(mul_8, unsqueeze_3)
	34	+ return add_5
	35	+
	36	+ def test_issue62(self):
	37	+ addmm_5 = torch.randn((2, 9216), device='npu:0', dtype=torch.float16)
	38	+ add = torch.randn((2, 4096, 1536), device='npu:0', dtype=torch.float16)
	39	+
	40	+ std_ret = self.op_func(addmm_5, add)
	41	+ compiled_func = torch.compile(self.op_func, backend="inductor")
	42	+ inductor_ret = compiled_func(addmm_5, add)
	43	+ self.assertEqual(std_ret, inductor_ret, atol=1e-2, rtol=1e-2)
	44	+
	45	+
	46	+if __name__ == "__main__":
	47	+ run_tests()

Atest/_inductor/test_issue70.py+22-0

@@ -0,0 +1,22 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class Test_issue70(TestUtils):
	8	+ def op_forward(self, x):
	9	+ return x.mean(-1)
	10	+
	11	+ def test_issue70(self):
	12	+ compiled_net = torch.compile(self.op_forward, backend="inductor")
	13	+
	14	+ arg = torch.randn((1, 1, 7168)).npu()
	15	+
	16	+ output = self.op_forward(arg)
	17	+ output1 = compiled_net(arg)
	18	+ self.assertEqual(output, output1, atol=1e-3, rtol=1e-3)
	19	+
	20	+
	21	+if __name__ == "__main__":
	22	+ run_tests()

Atest/_inductor/test_lazy_register.py+35-0

@@ -0,0 +1,35 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestLazyRegister(TestUtils):
	8	+ def test_compile_but_not_invoked(self):
	9	+
	10	+ def run(x, y):
	11	+ return x + y
	12	+
	13	+ run = torch.compile(run)
	14	+ self.assertFalse(torch_npu.utils._dynamo.is_inductor_npu_initialized())
	15	+
	16	+ def test_disale_register_inductor_npu(self):
	17	+ torch_npu.utils._dynamo.disable_register_inductor_npu()
	18	+
	19	+ def run(x, y):
	20	+ return x - y
	21	+
	22	+ run = torch.compile(run)
	23	+ x = torch.randn(10, 20).npu()
	24	+ y = torch.randn(10, 20).npu()
	25	+
	26	+ with self.assertRaisesRegex(Exception, "Device npu not supported"):
	27	+ _ = run(x, y)
	28	+
	29	+ self.assertFalse(torch_npu.utils._dynamo.is_inductor_npu_initialized())
	30	+
	31	+ torch_npu.utils._dynamo.enable_register_inductor_npu()
	32	+
	33	+
	34	+if __name__ == "__main__":
	35	+ run_tests()

Atest/_inductor/test_opensora_graph1.py+266-0

@@ -0,0 +1,266 @@
	1	+import os
	2	+import random
	3	+import numpy as np
	4	+
	5	+import torch
	6	+from torch import device
	7	+from torch.testing._internal.common_utils import run_tests
	8	+from testutils import TestUtils
	9	+import torch_npu
	10	+
	11	+device_npu = 'npu'
	12	+
	13	+
	14	+class TestModel(TestUtils):
	15	+ def test_opensora_cases_model_9_inference(self):
	16	+ def forward(primals_1: "f32[1, 9600, 2304]"):
	17	+ permute: "f32[9600, 1, 2304]" = torch.ops.aten.permute.default(primals_1, [1, 0, 2])
	18	+ return permute
	19	+ primals_2 = torch.randn((1, 9600, 2304), device=device_npu, dtype=torch.float32)
	20	+ ref = forward(primals_2)
	21	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	22	+ calc = forward_calc(primals_2)
	23	+ self.assertEqual(ref, calc, atol=1e-4, rtol=1e-4, equal_nan=True)
	24	+ primals_3 = torch.randn((1, 512, 2304), device=device_npu, dtype=torch.float32)
	25	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	26	+ calc = forward_calc(primals_3)
	27	+ ref = forward(primals_3)
	28	+ self.assertEqual(ref, calc, atol=1e-4, rtol=1e-4, equal_nan=True)
	29	+ primals_4 = torch.randn((9600, 1, 2304), device=device_npu, dtype=torch.float32)
	30	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	31	+ calc = forward_calc(primals_4)
	32	+ ref = forward(primals_4)
	33	+ self.assertEqual(ref, calc, atol=1e-4, rtol=1e-4, equal_nan=True)
	34	+
	35	+ def test_opensora_cases_model_11_inference(self):
	36	+ def forward(arg0_1: "f32[1, 1, 9600]", arg1_1: "f32[1, 1, 512]"):
	37	+ unsqueeze: "f32[1, 1, 1, 9600]" = torch.ops.aten.unsqueeze.default(arg0_1, 1)
	38	+ arg0_1 = None
	39	+ unsqueeze_1: "f32[1, 1, 1, 512]" = torch.ops.aten.unsqueeze.default(arg1_1, 1)
	40	+ arg1_1 = None
	41	+ constant_pad_nd: "f32[1, 1, 1, 9600]" = torch.ops.aten.constant_pad_nd.default(unsqueeze, [0, 0, 0, 0], -9980.0)
	42	+ unsqueeze = None
	43	+ view: "f32[1, 9600, 1]" = torch.ops.aten.view.default(constant_pad_nd, [1, 9600, 1])
	44	+ permute: "f32[1, 1, 9600]" = torch.ops.aten.permute.default(view, [2, 0, 1])
	45	+ view = None
	46	+ view_1: "f32[1, 1, 1, 9600]" = torch.ops.aten.view.default(permute, [1, 1, 1, 9600])
	47	+ permute = None
	48	+ view_2: "f32[1, 9600, 1, 1]" = torch.ops.aten.view.default(constant_pad_nd, [1, 9600, 1, 1])
	49	+ constant_pad_nd = None
	50	+ permute_1: "f32[1, 1, 9600, 1]" = torch.ops.aten.permute.default(view_2, [2, 0, 1, 3])
	51	+ view_2 = None
	52	+ view_3: "f32[1, 1, 1, 9600]" = torch.ops.aten.view.default(permute_1, [1, 1, 1, 9600])
	53	+ permute_1 = None
	54	+ repeat: "f32[1, 1, 1, 512]" = torch.ops.aten.repeat.default(unsqueeze_1, [1, 1, 1, 1])
	55	+ unsqueeze_1 = None
	56	+ npu_dtype_cast: "b8[1, 1, 1, 9600]" = torch.ops.npu.npu_dtype_cast.default(view_1, torch.bool)
	57	+ view_1 = None
	58	+ repeat_1: "b8[1, 1, 9600, 9600]" = torch.ops.aten.repeat.default(npu_dtype_cast, [1, 1, 9600, 1])
	59	+ npu_dtype_cast = None
	60	+ npu_dtype_cast_1: "b8[1, 1, 1, 9600]" = torch.ops.npu.npu_dtype_cast.default(view_3, torch.bool)
	61	+ view_3 = None
	62	+ repeat_2: "b8[1, 1, 9600, 9600]" = torch.ops.aten.repeat.default(npu_dtype_cast_1, [1, 1, 9600, 1])
	63	+ npu_dtype_cast_1 = None
	64	+ npu_dtype_cast_2: "b8[1, 1, 1, 512]" = torch.ops.npu.npu_dtype_cast.default(repeat, torch.bool)
	65	+ repeat = None
	66	+ repeat_3: "b8[1, 1, 9600, 512]" = torch.ops.aten.repeat.default(npu_dtype_cast_2, [1, 1, 9600, 1])
	67	+ npu_dtype_cast_2 = None
	68	+ return (repeat_1, repeat_3, repeat_2)
	69	+ arg0_1 = torch.rand((1, 1, 9600), device=device_npu, dtype=torch.float32)
	70	+ arg1_1 = torch.rand((1, 1, 512), device=device_npu, dtype=torch.float32)
	71	+ ref = forward(arg0_1, arg1_1)
	72	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	73	+ calc = forward_calc(arg0_1, arg1_1)
	74	+
	75	+ for r, c in zip(ref, calc):
	76	+ self.assertEqual(r, c, atol=1e-4, rtol=1e-4, equal_nan=True)
	77	+
	78	+
	79	+ def test_opensora_cases_model_14_backward(self):
	80	+ def forward(args):
	81	+ primals_5, getitem_3, rsqrt, add_2, view, permute_1, tangents_1 = args
	82	+ sub: "f32[1, 9600, 2304]" = torch.ops.aten.sub.Tensor(primals_5, getitem_3)
	83	+ mul: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(sub, rsqrt)
	84	+ view_2: "f32[9600, 32]" = torch.ops.aten.view.default(tangents_1, [9600, 32])
	85	+ mm: "f32[9600, 2304]" = torch.ops.aten.mm.default(view_2, permute_1)
	86	+ permute_2: "f32[32, 9600]" = torch.ops.aten.permute.default(view_2, [1, 0])
	87	+ mm_1: "f32[32, 2304]" = torch.ops.aten.mm.default(permute_2, view)
	88	+ permute_3: "f32[2304, 32]" = torch.ops.aten.permute.default(mm_1, [1, 0])
	89	+ sum_1: "f32[1, 32]" = torch.ops.aten.sum.dim_IntList(view_2, [0], True)
	90	+ view_3: "f32[32]" = torch.ops.aten.view.default(sum_1, [32])
	91	+ permute_4: "f32[32, 2304]" = torch.ops.aten.permute.default(permute_3, [1, 0])
	92	+ view_4: "f32[1, 9600, 2304]" = torch.ops.aten.view.default(mm, [1, 9600, 2304])
	93	+ sum_2: "f32[1, 1, 2304]" = torch.ops.aten.sum.dim_IntList(view_4, [1], True)
	94	+ mul_2: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(view_4, mul)
	95	+ mul_3: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(view_4, add_2)
	96	+ sum_3: "f32[1, 1, 2304]" = torch.ops.aten.sum.dim_IntList(mul_2, [1], True)
	97	+ mul_5: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(mul_3, 2304)
	98	+ sum_4: "f32[1, 9600, 1]" = torch.ops.aten.sum.dim_IntList(mul_3, [2], True)
	99	+ mul_6: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(mul_3, mul)
	100	+ sum_5: "f32[1, 9600, 1]" = torch.ops.aten.sum.dim_IntList(mul_6, [2], True)
	101	+ mul_7: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(mul, sum_5)
	102	+ sub_2: "f32[1, 9600, 2304]" = torch.ops.aten.sub.Tensor(mul_5, sum_4)
	103	+ sub_3: "f32[1, 9600, 2304]" = torch.ops.aten.sub.Tensor(sub_2, mul_7)
	104	+ div: "f32[1, 9600, 1]" = torch.ops.aten.div.Tensor(rsqrt, 2304)
	105	+ mul_8: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(div, sub_3)
	106	+ cat: "f32[1, 2, 2304]" = torch.ops.aten.cat.default([sum_2, sum_3], 1)
	107	+ sum_6: "f32[1, 1, 2304]" = torch.ops.aten.sum.dim_IntList(cat, [1], True)
	108	+ squeeze_1: "f32[1, 2304]" = torch.ops.aten.squeeze.dim(sum_6, 1)
	109	+ full_default: "f32[1, 2304]" = torch.ops.aten.full.default([1, 2304], 0, dtype=torch.float32, device='npu')
	110	+ slice_scatter: "f32[1, 2304]" = torch.ops.aten.slice_scatter.default(full_default, squeeze_1, 0, 0,
	111	+ 9223372036854775807)
	112	+ squeeze_2: "f32[2, 2304]" = torch.ops.aten.squeeze.dim(cat, 0)
	113	+ return [squeeze_2, permute_4, view_3, slice_scatter, mul_8]
	114	+ primals_5 = torch.randn((1, 9600, 2304), device=device_npu, dtype=torch.float32)
	115	+ getitem_3 = torch.randn((1, 9600, 1), device=device_npu, dtype=torch.float32)
	116	+ rsqrt = torch.randn((1, 9600, 1), device=device_npu, dtype=torch.float32)
	117	+ add_2 = torch.randn((1, 1, 2304), device=device_npu, dtype=torch.float32)
	118	+ view = torch.randn((9600, 2304), device=device_npu, dtype=torch.float32)
	119	+ permute_1 = torch.randn((32, 2304), device=device_npu, dtype=torch.float32)
	120	+ tangents_1 = torch.randn((1, 9600, 32), device=device_npu, dtype=torch.float32)
	121	+ args = (primals_5, getitem_3, rsqrt, add_2, view, permute_1, tangents_1)
	122	+ ref = forward(args)
	123	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	124	+ calc = forward_calc(args)
	125	+
	126	+ for r, c in zip(ref, calc):
	127	+ self.assertEqual(r, c, atol=1e-3, rtol=1e-3, equal_nan=True)
	128	+
	129	+
	130	+ def test_opensora_cases_model_14_forward(self):
	131	+ def forward(primals_1: "f32[2, 2304]", primals_2: "f32[32, 2304]", primals_3: "f32[32]",
	132	+ primals_4: "f32[1, 2304]", primals_5: "f32[1, 9600, 2304]"):
	133	+ unsqueeze: "f32[1, 2, 2304]" = torch.ops.aten.unsqueeze.default(primals_1, 0)
	134	+ primals_1 = None
	135	+ slice_1: "f32[1, 2304]" = torch.ops.aten.slice.Tensor(primals_4, 0, 0, 9223372036854775807)
	136	+ primals_4 = None
	137	+ unsqueeze_1: "f32[1, 1, 2304]" = torch.ops.aten.unsqueeze.default(slice_1, 1)
	138	+ slice_1 = None
	139	+ add: "f32[1, 2, 2304]" = torch.ops.aten.add.Tensor(unsqueeze, unsqueeze_1)
	140	+ unsqueeze = unsqueeze_1 = None
	141	+ split = torch.ops.aten.split.Tensor(add, 1, 1)
	142	+ add = None
	143	+ getitem: "f32[1, 1, 2304]" = split[0]
	144	+ getitem_1: "f32[1, 1, 2304]" = split[1]
	145	+ split = None
	146	+ var_mean = torch.ops.aten.var_mean.correction(primals_5, [2], correction=0, keepdim=True)
	147	+ getitem_2: "f32[1, 9600, 1]" = var_mean[0]
	148	+ getitem_3: "f32[1, 9600, 1]" = var_mean[1]
	149	+ var_mean = None
	150	+ add_1: "f32[1, 9600, 1]" = torch.ops.aten.add.Tensor(getitem_2, 1e-06)
	151	+ getitem_2 = None
	152	+ rsqrt: "f32[1, 9600, 1]" = torch.ops.aten.rsqrt.default(add_1)
	153	+ add_1 = None
	154	+ sub: "f32[1, 9600, 2304]" = torch.ops.aten.sub.Tensor(primals_5, getitem_3)
	155	+ mul: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(sub, rsqrt)
	156	+ sub = None
	157	+ add_2: "f32[1, 1, 2304]" = torch.ops.aten.add.Tensor(getitem_1, 1)
	158	+ getitem_1 = None
	159	+ mul_1: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(mul, add_2)
	160	+ mul = None
	161	+ add_3: "f32[1, 9600, 2304]" = torch.ops.aten.add.Tensor(mul_1, getitem)
	162	+ mul_1 = getitem = None
	163	+ view: "f32[9600, 2304]" = torch.ops.aten.view.default(add_3, [9600, 2304])
	164	+ add_3 = None
	165	+ permute: "f32[2304, 32]" = torch.ops.aten.permute.default(primals_2, [1, 0])
	166	+ primals_2 = None
	167	+ addmm: "f32[9600, 32]" = torch.ops.aten.addmm.default(primals_3, view, permute)
	168	+ primals_3 = None
	169	+ view_1: "f32[1, 9600, 32]" = torch.ops.aten.view.default(addmm, [1, 9600, 32])
	170	+ addmm = None
	171	+ # No stacktrace found for following nodes
	172	+ squeeze: "f32[1, 9600, 32]" = torch.ops.aten.squeeze.dim(view_1, 1)
	173	+ view_1 = None
	174	+ permute_1: "f32[32, 2304]" = torch.ops.aten.permute.default(permute, [1, 0])
	175	+ permute = None
	176	+ return [squeeze, primals_5, getitem_3, rsqrt, add_2, view, permute_1]
	177	+ primals_1 = torch.ones((2, 2304), device=device_npu, dtype=torch.float32)
	178	+ primals_2 = torch.ones((32, 2304), device=device_npu, dtype=torch.float32)
	179	+ primals_3 = torch.ones((32,), device=device_npu, dtype=torch.float32)
	180	+ primals_4 = torch.ones((1, 2304), device=device_npu, dtype=torch.float32)
	181	+ primals_5 = torch.ones((1, 9600, 2304), device=device_npu, dtype=torch.float32)
	182	+ ref = forward(primals_1, primals_2, primals_3, primals_4, primals_5)
	183	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	184	+ calc = forward_calc(primals_1, primals_2, primals_3, primals_4, primals_5)
	185	+ for r, c in zip(ref, calc):
	186	+ self.assertEqual(r, c, atol=1e-4, rtol=1e-4, equal_nan=True)
	187	+
	188	+
	189	+ def test_opensora_cases_model_15_forward(self):
	190	+ def forward(primals_1: "f32[1, 8, 30, 40, 1, 2, 2, 8]", primals_2: "i64[]", primals_3: "i64[]",
	191	+ primals_4: "i64[]"):
	192	+ permute: "f32[1, 8, 8, 1, 30, 2, 40, 2]" = torch.ops.aten.permute.default(primals_1, [0, 7, 1, 4, 2, 5, 3, 6])
	193	+ mul: "i64[]" = torch.ops.aten.mul.Tensor(primals_2, 1)
	194	+ mul_1: "i64[]" = torch.ops.aten.mul.Tensor(primals_3, 2)
	195	+ mul_2: "i64[]" = torch.ops.aten.mul.Tensor(primals_4, 2)
	196	+ return [permute, mul, mul_1, mul_2]
	197	+
	198	+ primals_1 = torch.randn((1, 8, 30, 40, 1, 2, 2, 8), device=device_npu, dtype=torch.float32)
	199	+ primals_2 = torch.tensor((1), device=device_npu, dtype=torch.int64)
	200	+ primals_3 = torch.tensor((1), device=device_npu, dtype=torch.int64)
	201	+ primals_4 = torch.tensor((1), device=device_npu, dtype=torch.int64)
	202	+ ref = forward(primals_1, primals_2, primals_3,
	203	+ primals_4)
	204	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	205	+ calc = forward_calc(primals_1, primals_2, primals_3,
	206	+ primals_4)
	207	+ for r, c in zip(ref, calc):
	208	+ self.assertEqual(r, c, atol=1e-4, rtol=1e-4, equal_nan=True)
	209	+
	210	+ def test_opensora_cases_model_16_forward(self):
	211	+ def forward(primals_1: "f32[2, 2304]", primals_2: "f32[32, 2304]", primals_3: "f32[32]", primals_4: "f32[1, 2304]", primals_5: "f32[1, 9600, 2304]"):
	212	+ unsqueeze: "f32[1, 2, 2304]" = torch.ops.aten.unsqueeze.default(primals_1, 0)
	213	+ slice_1: "f32[1, 2304]" = torch.ops.aten.slice.Tensor(primals_4, 0, 0, 9223372036854775807)
	214	+ unsqueeze_1: "f32[1, 1, 2304]" = torch.ops.aten.unsqueeze.default(slice_1, 1)
	215	+ add: "f32[1, 2, 2304]" = torch.ops.aten.add.Tensor(unsqueeze, unsqueeze_1)
	216	+ split = torch.ops.aten.split.Tensor(add, 1, 1)
	217	+ getitem: "f32[1, 1, 2304]" = split[0]
	218	+ getitem_1: "f32[1, 1, 2304]" = split[1]
	219	+ var_mean = torch.ops.aten.var_mean.correction(primals_5, [2], correction=0, keepdim=True)
	220	+ getitem_2: "f32[1, 9600, 1]" = var_mean[0]
	221	+ getitem_3: "f32[1, 9600, 1]" = var_mean[1]
	222	+ add_1: "f32[1, 9600, 1]" = torch.ops.aten.add.Tensor(getitem_2, 1e-06)
	223	+ rsqrt: "f32[1, 9600, 1]" = torch.ops.aten.rsqrt.default(add_1)
	224	+ sub: "f32[1, 9600, 2304]" = torch.ops.aten.sub.Tensor(primals_5, getitem_3)
	225	+ mul: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(sub, rsqrt)
	226	+ add_2: "f32[1, 1, 2304]" = torch.ops.aten.add.Tensor(getitem_1, 1)
	227	+ mul_1: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(mul, add_2)
	228	+ add_3: "f32[1, 9600, 2304]" = torch.ops.aten.add.Tensor(mul_1, getitem)
	229	+ view: "f32[9600, 2304]" = torch.ops.aten.view.default(add_3, [9600, 2304])
	230	+ permute: "f32[2304, 32]" = torch.ops.aten.permute.default(primals_2, [1, 0])
	231	+ addmm: "f32[9600, 32]" = torch.ops.aten.addmm.default(primals_3, view, permute)
	232	+ view_1: "f32[1, 9600, 32]" = torch.ops.aten.view.default(addmm, [1, 9600, 32])
	233	+ squeeze: "f32[1, 9600, 32]" = torch.ops.aten.squeeze.dim(view_1, 1)
	234	+ view_2: "f32[1, 8, 30, 40, 1, 2, 2, 8]" = torch.ops.aten.view.default(squeeze, [1, 8, 30, 40, 1, 2, 2, 8])
	235	+ permute_1: "f32[1, 8, 8, 1, 30, 2, 40, 2]" = torch.ops.aten.permute.default(view_2, [0, 7, 1, 4, 2, 5, 3, 6])
	236	+ clone: "f32[1, 8, 8, 1, 30, 2, 40, 2]" = torch.ops.aten.clone.default(permute_1)
	237	+ clone_1: "f32[1, 8, 8, 1, 30, 2, 40, 2]" = torch.ops.aten.clone.default(clone, memory_format=torch.contiguous_format)
	238	+ view_3: "f32[1, 8, 8, 60, 80]" = torch.ops.aten.view.default(clone_1, [1, 8, 8, 60, 80])
	239	+ permute_3: "f32[32, 2304]" = torch.ops.aten.permute.default(permute, [1, 0])
	240	+ return [view_3, primals_5, getitem_3, rsqrt, add_2, view, permute_3]
	241	+
	242	+ def seed_all(seed=1234, mode=False):
	243	+ random.seed(seed)
	244	+ os.environ['PYTHONHASHSEED'] = str(seed)
	245	+ np.random.seed(seed)
	246	+ torch.manual_seed(seed)
	247	+ torch.use_deterministic_algorithms(mode)
	248	+ torch_npu.npu.manual_seed_all(seed)
	249	+ torch_npu.npu.manual_seed(seed)
	250	+
	251	+ seed_all(True)
	252	+ primals_1 = torch.randn((2, 2304), device=device_npu, dtype=torch.float32)
	253	+ primals_2 = torch.randn((32, 2304), device=device_npu, dtype=torch.float32)
	254	+ primals_3 = torch.randn((32,), device=device_npu, dtype=torch.float32)
	255	+ primals_4 = torch.randn((1, 2304), device=device_npu, dtype=torch.float32)
	256	+ primals_5 = torch.randn((1, 9600, 2304), device=device_npu, dtype=torch.float32)
	257	+
	258	+ ref = forward(primals_1, primals_2, primals_3, primals_4, primals_5)
	259	+ forward_calc = torch.compile(forward, backend="inductor", dynamic=False)
	260	+ calc = forward_calc(primals_1, primals_2, primals_3, primals_4, primals_5)
	261	+ for r, c in zip(ref, calc):
	262	+ self.assertEqual(r, c, atol=1e-3, rtol=1e-3, equal_nan=True)
	263	+
	264	+
	265	+if __name__ == "__main__":
	266	+ run_tests()

Atest/_inductor/test_permute.py+39-0

@@ -0,0 +1,39 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestPermute(TestUtils):
	8	+ _permute_dims = [
	9	+ (0, 1, 2, 3), (0, 1, 3, 2), (0, 2, 1, 3), (0, 2, 3, 1),
	10	+ (0, 3, 1, 2), (0, 3, 2, 1), (1, 0, 2, 3), (1, 0, 3, 2),
	11	+ (1, 2, 0, 3), (1, 2, 3, 0), (1, 3, 0, 2), (1, 3, 2, 0),
	12	+ (2, 0, 1, 3), (2, 0, 3, 1), (2, 1, 0, 3), (2, 1, 3, 0),
	13	+ (2, 3, 0, 1), (2, 3, 1, 0), (3, 0, 1, 2), (3, 0, 2, 1),
	14	+ (3, 1, 0, 2), (3, 1, 2, 0), (3, 2, 0, 1), (3, 2, 1, 0),
	15	+ ]
	16	+
	17	+ def op_calc(self, a, b, dim):
	18	+ a = a.permute(dim)
	19	+ b = b.permute(dim)
	20	+ y = a + b
	21	+ return y
	22	+
	23	+ @parametrize('shape', [(8, 8, 512, 128)])
	24	+ @parametrize('dtype', ['float32', 'int32', 'float16', 'bfloat16', 'int64'])
	25	+ def test_view_cases(self, shape, dtype):
	26	+ a = self._generate_tensor(shape, dtype)
	27	+ b = self._generate_tensor(shape, dtype)
	28	+
	29	+ for dim in self._permute_dims:
	30	+ std_permute = self.op_calc(a, b, dim)
	31	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	32	+ inductor_permute = compiled_op_calc(a, b, dim)
	33	+
	34	+ self.assertEqual(std_permute, inductor_permute, atol=1e-3, rtol=1e-3)
	35	+
	36	+instantiate_parametrized_tests(TestPermute)
	37	+
	38	+if __name__ == "__main__":
	39	+ run_tests()

Atest/_inductor/test_reduction_brocast_add.py+30-0

@@ -0,0 +1,30 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSumAdd(TestUtils):
	8	+ def foo(self, a, b, dim, shape):
	9	+ y = a + b
	10	+ y = y.sum(dim)
	11	+ y = y.unsqueeze(dim)
	12	+ y = y.broadcast_to(shape) + b
	13	+ return y
	14	+
	15	+ # case：change shapes
	16	+ @parametrize('shape', [(9, 9, 31, 63)])
	17	+ @parametrize('dim', [0, 1, 2])
	18	+ @parametrize('dtype', ['float32'])
	19	+ def test_reduction_cases_shapes1(self, shape, dim, dtype):
	20	+ a, b = [torch.randn(shape, requires_grad=False, dtype=torch.float32, device="npu") for _ in range(2)]
	21	+ r1 = self.foo(a, b, dim, shape)
	22	+ func = torch.compile(self.foo, backend="inductor", dynamic=False)
	23	+ r = func(a, b, dim, shape)
	24	+ self.assertEqual(r, r1, atol=1e-3, rtol=1e-3)
	25	+
	26	+
	27	+instantiate_parametrized_tests(TestSumAdd)
	28	+
	29	+if __name__ == "__main__":
	30	+ run_tests()

Atest/_inductor/test_relu.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestRelu(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.relu(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+
	22	+ self.assertEqual(std_result, inductor_result)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestRelu)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_renorm.py+27-0

@@ -0,0 +1,27 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestRenorm(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return torch.renorm(input_element, p=2, dim=dim, maxnorm=5)
	10	+
	11	+ # case：change shapes
	12	+ @parametrize('shape', [(32, 64)])
	13	+ @parametrize('dim', [-1])
	14	+ @parametrize('dtype', ['float32'])
	15	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	16	+ input_element = self._generate_tensor(shape, dtype)
	17	+ std_ret = self.op_calc(input_element, dim)
	18	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	19	+ inductor_ret = compiled_op_calc(input_element, dim)
	20	+
	21	+ self.assertEqual(std_ret, inductor_ret, equal_nan=True)
	22	+
	23	+
	24	+instantiate_parametrized_tests(TestRenorm)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_repeat.py+29-0

@@ -0,0 +1,29 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestRepeat(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return input_element.repeat(dim)
	10	+
	11	+ # case：change shapes
	12	+ @parametrize('shape', [(16, 128, 64)])
	13	+ @parametrize('dim', [(1, 1, 2), (1, 2, 1), (2, 1, 1)])
	14	+ @parametrize('dtype', ['float32'])
	15	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	16	+ input_element = self._generate_tensor(shape, dtype)
	17	+
	18	+ std_ret = self.op_calc(input_element, dim)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	21	+ inductor_ret = compiled_op_calc(input_element, dim)
	22	+
	23	+ self.assertEqual(std_ret, inductor_ret, atol=1e-1, rtol=1e-1)
	24	+
	25	+
	26	+instantiate_parametrized_tests(TestRepeat)
	27	+
	28	+if __name__ == "__main__":
	29	+ run_tests()

Atest/_inductor/test_reshape.py+35-0

@@ -0,0 +1,35 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+
	8	+
	9	+class TestReshape(TestUtils):
	10	+ B, N, S, D = (1, 12, 256, 8)
	11	+
	12	+ def op_calc(self, a, b):
	13	+ a = a.reshape(self.S, self.B, self.N * self.D)
	14	+ b = b.reshape(self.S, self.B, self.N * self.D)
	15	+ y = a + b
	16	+ return y
	17	+
	18	+ @parametrize('shape', [(1, 12, 256, 8)])
	19	+ @parametrize('dtype', ['float32', 'int32', 'float16', 'bfloat16', 'int64'])
	20	+ def test_view_cases(self, shape, dtype):
	21	+ a = self._generate_tensor(shape, dtype)
	22	+ b = self._generate_tensor(shape, dtype)
	23	+
	24	+ std_reshape = self.op_calc(a, b)
	25	+
	26	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	27	+ inductor_reshape = compiled_op_calc(a, b)
	28	+
	29	+ self.assertEqual(std_reshape, inductor_reshape, atol=1e-3, rtol=1e-3)
	30	+
	31	+
	32	+instantiate_parametrized_tests(TestReshape)
	33	+
	34	+if __name__ == "__main__":
	35	+ run_tests()

Atest/_inductor/test_rsqrt.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestRsqrt(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.rsqrt(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype, 1)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+
	22	+ self.assertEqual(std_result, inductor_result, atol=1e-1, rtol=1e-1)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestRsqrt)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_slice.py+50-0

@@ -0,0 +1,50 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSlice(TestUtils):
	8	+ def op_calc(self, a, b, dim, step):
	9	+ if dim == 0:
	10	+ target = a.shape[0]
	11	+ end = target // step
	12	+ a = a[:end:, ::, ::, ::]
	13	+ b = b[:end:, ::, ::, ::]
	14	+ elif dim == 1:
	15	+ target = a.shape[1]
	16	+ end = target // step
	17	+ a = a[::, :end:, ::, ::]
	18	+ b = b[::, :end:, ::, ::]
	19	+ elif dim == 2:
	20	+ target = a.shape[2]
	21	+ end = target // step
	22	+ a = a[::, ::, :end:, ::]
	23	+ b = b[::, ::, :end:, ::]
	24	+ elif dim == 3:
	25	+ target = a.shape[3]
	26	+ end = target // step
	27	+ a = a[::, ::, ::, :end:]
	28	+ b = b[::, ::, ::, :end:]
	29	+ y = a + b
	30	+ return y
	31	+
	32	+ @parametrize('shape', [(8, 8, 256, 128)])
	33	+ @parametrize('dtype', ['float32', 'int32', 'float16', 'bfloat16', 'int64'])
	34	+ def test_view_cases(self, shape, dtype):
	35	+ a = self._generate_tensor(shape, dtype)
	36	+ b = self._generate_tensor(shape, dtype)
	37	+
	38	+ for dim in [3, 2, 1, 0]:
	39	+ std_slice = self.op_calc(a, b, dim, min(shape) // 2)
	40	+
	41	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	42	+ inductor_slice = compiled_op_calc(a, b, dim, min(shape) // 2)
	43	+
	44	+ self.assertEqual(std_slice, inductor_slice, atol=1e-3, rtol=1e-3)
	45	+
	46	+
	47	+instantiate_parametrized_tests(TestSlice)
	48	+
	49	+if __name__ == "__main__":
	50	+ run_tests()

Atest/_inductor/test_split_loop.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSplitLoop(TestUtils):
	8	+ def op_calc(self, a, b):
	9	+ return torch.nn.functional.gelu(a + b)
	10	+
	11	+ @parametrize('shape', [(8, 86, 1152), (61, 89, 157), (7, 89, 971)])
	12	+ @parametrize('dtype', ['float32'])
	13	+ def test_split_loop(self, shape, dtype):
	14	+
	15	+ a = self._generate_tensor(shape, dtype)
	16	+ b = self._generate_tensor((shape[0], 1, shape[2]), dtype)
	17	+
	18	+ std_ = self.op_calc(a, b)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	21	+ inductor_ = compiled_op_calc(a, b)
	22	+ self.assertEqual(std_, inductor_, atol=1e-3, rtol=1e-3)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestSplitLoop)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_sqrt.py+27-0

@@ -0,0 +1,27 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSqrt(TestUtils):
	8	+ def op_calc(self, first_element):
	9	+ result = torch.sqrt(first_element)
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype, 1)
	16	+
	17	+ std_result = self.op_calc(first_element)
	18	+
	19	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	20	+ inductor_result = compiled_op_calc(first_element)
	21	+ self.assertEqual(std_result, inductor_result, atol=1e-1, rtol=1e-1, equal_nan=True)
	22	+
	23	+
	24	+instantiate_parametrized_tests(TestSqrt)
	25	+
	26	+if __name__ == "__main__":
	27	+ run_tests()

Atest/_inductor/test_sub.py+28-0

@@ -0,0 +1,28 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSub(TestUtils):
	8	+ def op_calc(self, first_element, second_element):
	9	+ result = first_element - second_element
	10	+ return result
	11	+
	12	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	13	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
	14	+ def test_pointwise_cases(self, shape, dtype):
	15	+ first_element = self._generate_tensor(shape, dtype)
	16	+ second_element = self._generate_tensor(shape, dtype)
	17	+
	18	+ std_sub = self.op_calc(first_element, second_element)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_sum = compiled_op_calc(first_element, second_element)
	22	+ self.assertEqual(std_sub, inductor_sum)
	23	+
	24	+
	25	+instantiate_parametrized_tests(TestSub)
	26	+
	27	+if __name__ == "__main__":
	28	+ run_tests()

Atest/_inductor/test_sum.py+55-0

@@ -0,0 +1,55 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSum(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return torch.sum(input_element, dim)
	10	+ # 规约轴和非规约轴对齐用例 float32 XBLOCK_SUB>=8:shape=(8,32)
	11	+ # non-persistent reduction 用例规约轴>1024:shape=(8,8,8,2048) dim=-1
	12	+ _reduction_extest_shape4d_all = [(8, 32), (8, 8, 8, 2048)]
	13	+ _reduction_extest_dim4d_low = [-1]
	14	+ _reduction_extest_dim4d_all = [0, 1, 2]
	15	+
	16	+ @parametrize('shape', _reduction_extest_shape4d_all)
	17	+ @parametrize('dim', _reduction_extest_dim4d_low)
	18	+ @parametrize('dtype', ['float32'])
	19	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	20	+ input_element = self._generate_tensor(shape, dtype)
	21	+ std_sum = self.op_calc(input_element, dim)
	22	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	23	+ inductor_sum_tmp = compiled_op_calc(input_element, dim)
	24	+ if dtype == 'int32' or dtype == 'int64':
	25	+ # inductor return float32,need to change int64 for assert
	26	+ inductor_sum = inductor_sum_tmp.long()
	27	+ elif dtype == 'float16':
	28	+ # inductor return float32,need to change float16 for assert
	29	+ inductor_sum = inductor_sum_tmp.half()
	30	+ elif dtype == 'bfloat16':
	31	+ # inductor return float32,need to change float32 for assert
	32	+ std_sum = std_sum.float()
	33	+ inductor_sum = inductor_sum_tmp
	34	+ else:
	35	+ inductor_sum = inductor_sum_tmp
	36	+
	37	+ self.assertEqual(std_sum, inductor_sum, atol=1e-1, rtol=1e-1)
	38	+
	39	+ @parametrize('shape', [(32, 16, 64, 128)])
	40	+ @parametrize('dim', _reduction_extest_dim4d_all)
	41	+ @parametrize('dtype', ['float32'])
	42	+ def test_reduction_cases_dims(self, shape, dim, dtype):
	43	+
	44	+ input_element = self._generate_tensor(shape, dtype)
	45	+ std_sum = self.op_calc(input_element, dim)
	46	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	47	+ inductor_sum = compiled_op_calc(input_element, dim)
	48	+
	49	+ self.assertEqual(std_sum, inductor_sum, atol=1e-1, rtol=1e-1)
	50	+
	51	+
	52	+instantiate_parametrized_tests(TestSum)
	53	+
	54	+if __name__ == "__main__":
	55	+ run_tests()

Atest/_inductor/test_sum_add.py+37-0

@@ -0,0 +1,37 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestSumAdd(TestUtils):
	8	+ def op_calc(self, input_element, dim, input_element2):
	9	+ tmp = torch.sum(input_element, dim)
	10	+ return tmp + input_element2
	11	+
	12	+ @parametrize('shape', [(32, 64, 128, 2048)])
	13	+ @parametrize('dim', [0, 1, 2, 3])
	14	+ @parametrize('dtype', ['float32'])
	15	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	16	+ input_element = self._generate_tensor(shape, dtype)
	17	+ if dim == -1 or dim == 3:
	18	+ input_element2 = torch.full(size=(32, 64, 128), fill_value=1000.0, dtype=torch.float32, device=torch.device("npu"))
	19	+ elif dim == 2:
	20	+ input_element2 = torch.full(size=(32, 64, 2048), fill_value=1000.0, dtype=torch.float32, device=torch.device("npu"))
	21	+ elif dim == 1:
	22	+ input_element2 = torch.full(size=(32, 128, 2048), fill_value=1000.0, dtype=torch.float32, device=torch.device("npu"))
	23	+ else:
	24	+ input_element2 = torch.full(size=(64, 128, 2048), fill_value=1000.0, dtype=torch.float32, device=torch.device("npu"))
	25	+
	26	+ std_sum = self.op_calc(input_element, dim, input_element2)
	27	+
	28	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	29	+ inductor_sum = compiled_op_calc(input_element, dim, input_element2)
	30	+
	31	+ self.assertEqual(std_sum, inductor_sum, atol=1e-1, rtol=1e-1)
	32	+
	33	+
	34	+instantiate_parametrized_tests(TestSumAdd)
	35	+
	36	+if __name__ == "__main__":
	37	+ run_tests()

Atest/_inductor/test_triton.py+36-0

@@ -0,0 +1,36 @@
	1	+# Owner(s): ["module: tests"]
	2	+
	3	+import unittest
	4	+import torch
	5	+from torch.testing._internal.common_utils import run_tests, TestCase, load_tests
	6	+from torch.utils._triton import has_triton_package, has_triton, has_triton_tma, has_triton_tma_device
	7	+import torch_npu
	8	+import torch_npu.testing
	9	+
	10	+# load_tests from torch.testing._internal.common_utils is used to automatically filter tests for
	11	+# sharding on sandcastle. This line silences flake warnings
	12	+load_tests = load_tests
	13	+
	14	+
	15	+class TestHasTriton(TestCase):
	16	+
	17	+ def test_has_triton(self):
	18	+ if not has_triton_package():
	19	+ # no triton library found, skip test_has_triton
	20	+ return
	21	+
	22	+ self.assertFalse(has_triton())
	23	+ self.assertFalse(has_triton_tma())
	24	+ self.assertFalse(has_triton_tma_device())
	25	+
	26	+ from torch_npu.contrib import transfer_to_npu
	27	+
	28	+ self.assertFalse(has_triton())
	29	+ self.assertFalse(has_triton_tma())
	30	+ self.assertFalse(has_triton_tma_device())
	31	+
	32	+
	33	+
	34	+
	35	+if __name__ == "__main__":
	36	+ run_tests()

Atest/_inductor/test_var.py+26-0

@@ -0,0 +1,26 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestVar(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return torch.var(input_element, dim)
	10	+
	11	+ # case：change shapes
	12	+ @parametrize('shape', [(8, 64, 128)])
	13	+ @parametrize('dim', [0, 1, 2])
	14	+ @parametrize('dtype', ['float16'])
	15	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	16	+ input_element = self._generate_tensor(shape, dtype)
	17	+ std_ret = self.op_calc(input_element, dim)
	18	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	19	+ inductor_ret = compiled_op_calc(input_element, dim)
	20	+ self.assertEqual(std_ret, inductor_ret, atol=1e-1, rtol=True, equal_nan=True)
	21	+
	22	+
	23	+instantiate_parametrized_tests(TestVar)
	24	+
	25	+if __name__ == "__main__":
	26	+ run_tests()

Atest/_inductor/test_var_mean.py+31-0

@@ -0,0 +1,31 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestVarMean(TestUtils):
	8	+ def op_calc(self, input_element, dim):
	9	+ return torch.var_mean(input_element, dim)
	10	+
	11	+ # case：The shape must not be too large
	12	+ @parametrize('shape', [(8, 64, 128)])
	13	+ @parametrize('dim', [0, 1, 2, (0, 2), (0, 1)])
	14	+ @parametrize('dtype', ['float32'])
	15	+ def test_reduction_cases_shapes(self, shape, dim, dtype):
	16	+
	17	+ input_element = self._generate_tensor(shape, dtype)
	18	+
	19	+ std_var, std_mean = self.op_calc(input_element, dim)
	20	+
	21	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor", dynamic=False)
	22	+ inductor_var, inductor_mean = compiled_op_calc(input_element, dim)
	23	+
	24	+ self.assertEqual(std_var, inductor_var, atol=1e-1, rtol=1e-1, equal_nan=True)
	25	+ self.assertEqual(std_mean, inductor_mean, atol=1e-1, rtol=1e-1, equal_nan=True)
	26	+
	27	+
	28	+instantiate_parametrized_tests(TestVarMean)
	29	+
	30	+if __name__ == "__main__":
	31	+ run_tests()

Atest/_inductor/test_var_mean_add_mul.py+48-0

@@ -0,0 +1,48 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestReduction(TestUtils):
	8	+ def forward(self, add: "f32[1, 2, 2304]", primals_2: "f32[32, 2304]", primals_5: "f32[1, 9600, 2304]"):
	9	+ split = torch.ops.aten.split.Tensor(add, 1, 1)
	10	+ getitem: "f32[1, 1, 2304]" = split[0]
	11	+ getitem_1: "f32[1, 1, 2304]" = split[1]
	12	+
	13	+ var_mean = torch.ops.aten.var_mean.correction(primals_5, [2], correction=0, keepdim=True)
	14	+ getitem_2: "f32[1, 9600, 1]" = var_mean[0]
	15	+ getitem_3: "f32[1, 9600, 1]" = var_mean[1]
	16	+ add_1: "f32[1, 9600, 1]" = torch.ops.aten.add.Tensor(getitem_2, 1e-06)
	17	+ rsqrt: "f32[1, 9600, 1]" = torch.ops.aten.rsqrt.default(add_1)
	18	+ sub: "f32[1, 9600, 2304]" = torch.ops.aten.sub.Tensor(primals_5, getitem_3)
	19	+ mul: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(sub, rsqrt)
	20	+
	21	+ add_2: "f32[1, 1, 2304]" = torch.ops.aten.add.Tensor(getitem_1, 1)
	22	+ mul_1: "f32[1, 9600, 2304]" = torch.ops.aten.mul.Tensor(mul, add_2)
	23	+ add_3: "f32[1, 9600, 2304]" = torch.ops.aten.add.Tensor(mul_1, getitem)
	24	+
	25	+ view: "f32[9600, 2304]" = torch.ops.aten.view.default(add_3, [9600, 2304])
	26	+ return [None, primals_5, getitem_3, rsqrt, add_2, view, primals_2]
	27	+
	28	+ def test_reduction_cases_shapes(self):
	29	+ device = 'npu'
	30	+ primals_2: "f32[32, 2304]" = torch.randn((32, 2304), device=device, dtype=torch.float32)
	31	+ primals_5: "f32[1, 9600, 2304]" = torch.randn((1, 9600, 2304), device=device, dtype=torch.float32)
	32	+ add: "f32[1, 2, 2304]" = torch.randn((1, 2, 2304), device=device, dtype=torch.float32)
	33	+
	34	+ _, primals_5_ref, getitem_3_ref, rsqrt_ref, add_2_ref, view_ref, primals_2_ref = self.forward(add, primals_2, primals_5)
	35	+
	36	+ self.forward = torch.compile(self.forward, backend="inductor", dynamic=False)
	37	+ _, primals_5, getitem_3, rsqrt, add_2, view, primals_2 = self.forward(add, primals_2, primals_5)
	38	+
	39	+ self.assertEqual(primals_5_ref, primals_5, atol=1e-3, rtol=1e-3, equal_nan=True)
	40	+ self.assertEqual(getitem_3_ref, getitem_3, atol=1e-3, rtol=1e-3, equal_nan=True)
	41	+ self.assertEqual(rsqrt_ref, rsqrt, atol=1e-3, rtol=1e-3, equal_nan=True)
	42	+ self.assertEqual(add_2_ref, add_2, atol=1e-3, rtol=1e-3, equal_nan=True)
	43	+ self.assertEqual(view_ref, view, atol=1e-3, rtol=1e-3, equal_nan=True)
	44	+ self.assertEqual(primals_2_ref, primals_2, atol=1e-3, rtol=1e-3, equal_nan=True)
	45	+
	46	+
	47	+if __name__ == "__main__":
	48	+ run_tests()

Atest/_inductor/test_where.py+29-0

@@ -0,0 +1,29 @@
	1	+import torch
	2	+from torch.testing._internal.common_utils import run_tests, parametrize, instantiate_parametrized_tests
	3	+from testutils import TestUtils
	4	+import torch_npu
	5	+
	6	+
	7	+class TestWhere(TestUtils):
	8	+ def op_calc(self, condition, first_element, second_element):
	9	+ return torch.where(condition, first_element, second_element)
	10	+
	11	+ @parametrize('shape', TestUtils._pointwise_demo_shapes)
	12	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32'])
	13	+ def test_pointwise_cases(self, shape, dtype):
	14	+ first_element = self._generate_tensor(shape, dtype)
	15	+ second_element = self._generate_tensor(shape, dtype)
	16	+ condition = self._generate_tensor(shape, 'bool')
	17	+
	18	+ std_result = self.op_calc(condition, first_element, second_element)
	19	+
	20	+ compiled_op_calc = torch.compile(self.op_calc, backend="inductor")
	21	+ inductor_result = compiled_op_calc(condition, first_element, second_element)
	22	+
	23	+ self.assertEqual(std_result, inductor_result)
	24	+
	25	+
	26	+instantiate_parametrized_tests(TestWhere)
	27	+
	28	+if __name__ == "__main__":
	29	+ run_tests()

Atest/_inductor/testutils.py+36-0

@@ -0,0 +1,36 @@
	1	+from collections.abc import Sequence
	2	+import os
	3	+import time
	4	+import numpy as np
	5	+import torch
	6	+from torch.testing._internal.common_utils import TestCase
	7	+import torch_npu
	8	+
	9	+
	10	+class TestUtils(TestCase):
	11	+ _pointwise_test_shape2d = [(4096, 256), (1024, 32), (8, 2048), (8, 4096)] # (8, 4), (8, 8), not supported
	12	+ _pointwise_test_shape3d = [(8, 8, 4), (8, 8, 8), (8, 8, 2048), (8, 8, 4096)]
	13	+ _pointwise_test_shape4d = [(128, 128, 4096, 4), (128, 128, 4096, 8),
	14	+ (32, 32, 1024, 1024)] # 1281284096*2048 is too big(512G)
	15	+ _pointwise_test_shapes = _pointwise_test_shape2d + _pointwise_test_shape3d + _pointwise_test_shape4d
	16	+
	17	+ _pointwise_demo_shapes = [(1024, 32), (8, 16, 256, 32)]
	18	+ _reduction_extest_shape4d = [(8, 8, 8, 16384), (8, 8, 16384, 8), (8, 16384, 8, 8), (16384, 8, 8, 8)]
	19	+ _reduction_extest_dim4d = [-1, -2, 1, 0]
	20	+ _reduction_extest_SDbinding = list(zip(_reduction_extest_shape4d, _reduction_extest_dim4d))
	21	+
	22	+ _test_dtypes = ['float32', 'int32', 'float16', 'bfloat16', 'int64']
	23	+
	24	+ @staticmethod
	25	+ def _generate_tensor(shape, dtype, floatPOSIFLAG=0):
	26	+ if dtype == 'float32' or dtype == 'float16' or dtype == 'bfloat16':
	27	+ if floatPOSIFLAG:
	28	+ return 1000 * torch.rand(size=shape, dtype=eval('torch.' + dtype), device=torch.device("npu"))
	29	+ else:
	30	+ return torch.randn(size=shape, dtype=eval('torch.' + dtype), device=torch.device("npu")) * 2000
	31	+ elif dtype == 'int32' or dtype == 'int64':
	32	+ return torch.randint(low=0, high=2000, size=shape, dtype=eval('torch.' + dtype), device=torch.device("npu"))
	33	+ elif dtype == 'bool':
	34	+ return torch.randint(low=0, high=2, size=shape, device=torch.device("npu")).bool()
	35	+ else:
	36	+ raise ValueError('Invalid parameter \"dtype\" is found : {}'.format(dtype))

Mtorch_npu/_inductor/__init__.py+9-2

	@@ -21,6 +21,7 @@ from .npu_choices import should_use_persistent_reduction
21	from .npu_device import NewNPUDeviceOpOverrides	21	from .npu_device import NewNPUDeviceOpOverrides
22	from .runtime import _load_cached_autotuning	22	from .runtime import _load_cached_autotuning
23	from .utils import get_current_raw_stream, patch_is_gpu, patch_has_triton	23	from .utils import get_current_raw_stream, patch_is_gpu, patch_has_triton
		24	+from .codecache import patch_aot_code_compiler_compile, patch_cache_base_get_system
24		25
25	set_compile_threads()	26	set_compile_threads()
26		27
	@@ -54,14 +55,16 @@ def patch_torch_for_aoti():
54	from .utils import patch_is_same_tensor	55	from .utils import patch_is_same_tensor
55	from .fx_passes.joint_graph import patch_constant_fold_uniform_value	56	from .fx_passes.joint_graph import patch_constant_fold_uniform_value
56	from .ir import patch_fallback_kernel_codegen	57	from .ir import patch_fallback_kernel_codegen
57	- from .codecache import patch_aot_code_compiler_compile	58	+
58	patch_codegen_with_cpp_wrapper()	59	patch_codegen_with_cpp_wrapper()
59	patch_get_cpp_torch_device_options()	60	patch_get_cpp_torch_device_options()
60	patch_device_to_aten()	61	patch_device_to_aten()
61	patch_is_same_tensor()	62	patch_is_same_tensor()
62	patch_constant_fold_uniform_value()	63	patch_constant_fold_uniform_value()
63	patch_fallback_kernel_codegen()	64	patch_fallback_kernel_codegen()
		DW D duck72162025年8月15日这个patch的位置参考2.6放，这个不是aoti的patch W wangl12592025年8月29日已修改
64	- patch_aot_code_compiler_compile()	65	+
		66	+ patch_aot_code_compiler_compile()
		67	+
65		68
66		69
67	if os.environ.get("DISABLE_AOTI_PATCH", "0") != "1":	70	if os.environ.get("DISABLE_AOTI_PATCH", "0") != "1":
	@@ -99,5 +102,9 @@ InductorChoices.should_use_persistent_reduction = should_use_persistent_reductio
99	autotune_cache._load_cached_autotuning = _load_cached_autotuning	102	autotune_cache._load_cached_autotuning = _load_cached_autotuning
100		103
101	register_fa_pass()	104	register_fa_pass()
		105	+patch_cache_base_get_system()
102	patch_is_gpu()	106	patch_is_gpu()
103	patch_has_triton()	107	patch_has_triton()
		108	+
		109	+
		110	+

Mtorch_npu/_inductor/codecache.py+50-2

	@@ -1,5 +1,7 @@
1	import os	1	import os
2	import contextlib	2	import contextlib
		3	+import hashlib
		4	+import json
3	from typing import (	5	from typing import (
4	Any,	6	Any,
5	Callable,	7	Callable,
	@@ -18,9 +20,9 @@ from typing import (
18		20
19	import torch	21	import torch
20	from torch._inductor import config	22	from torch._inductor import config
21	-from torch._inductor.codecache import get_lock_dir, LOCK_TIMEOUT	23	+from torch._inductor.codecache import CacheBase, get_lock_dir, LOCK_TIMEOUT
22	from torch._inductor.graph import GraphLowering	24	from torch._inductor.graph import GraphLowering
23	-	25	+import torch_npu
24	from torch_npu.utils._error_code import ErrCode, pta_error	26	from torch_npu.utils._error_code import ErrCode, pta_error
25		27
26	empty_json = "{}"	28	empty_json = "{}"
	@@ -35,6 +37,52 @@ def lock_context(key):
35	yield	37	yield
36		38
37		39
		40	+
		41	+def patch_cache_base_get_system():
		42	+ # patch function CacheBase.get_system with get_system_npu, add logic to support CANN
		43	+ @staticmethod
		44	+ def get_system():
		45	+ try:
		46	+ from triton.compiler.compiler import triton_key
		47	+
		48	+ # Use triton_key instead of triton.__version__ as the version
		49	+ # is not updated with each code change
		50	+ triton_version = triton_key()
		51	+ except ModuleNotFoundError:
		52	+ triton_version = None
		53	+
		54	+ try:
		55	+ system: Dict[str, Any] = {
		56	+ "device": {"name": None},
		57	+ "version": {
		58	+ "triton": triton_version,
		59	+ },
		60	+ }
		61	+ device_properties = torch_npu.npu.get_device_properties(
		62	+ torch_npu.npu.current_device()
		63	+ )
		64	+ if torch.version.cann is not None:
		65	+ system["device"]["name"] = device_properties.name
		66	+ system["version"]["cann"] = torch.version.cann
		67	+ elif torch.version.cuda is not None:
		68	+ system["device"]["name"] = device_properties.name
		69	+ system["version"]["cuda"] = torch.version.cuda
		70	+ else:
		71	+ system["device"]["name"] = device_properties.gcnArchName
		72	+ system["version"]["hip"] = torch.version.hip
		73	+ except (AssertionError, RuntimeError):
		74	+ # If deivce is not installed, none of the above config is relevant.
		75	+ system = {}
		76	+
		77	+ system["hash"] = hashlib.sha256(
		78	+ json.dumps(system, sort_keys=True).encode("utf-8")
		79	+ ).hexdigest()
		80	+
		81	+ return system
		82	+
		83	+ CacheBase.get_system = get_system
		84	+
		85	+
38	def patch_aot_code_compiler_compile():	86	def patch_aot_code_compiler_compile():
39	# In v2.6.0, aoti has bug when init oss_proxy_executor with default op_json,	87	# In v2.6.0, aoti has bug when init oss_proxy_executor with default op_json,
40	# which could not be skipped, so here we try to create a new npu op_json,	88	# which could not be skipped, so here we try to create a new npu op_json,

Mtorch_npu/_inductor/lowering.py+5-1

	@@ -33,7 +33,7 @@ from torch._inductor.lowering import (
33	add_layout_constraint	33	add_layout_constraint
34	)	34	)
35	import torch_npu	35	import torch_npu
36	-from torch_npu import npu_dtype_cast	36	+from torch_npu import npu_dtype_cast, _npu_dtype_cast
37	from .lowering_op_list import GENERATE_LIST, GENERATE_LIST2, FALLBACK_LIST, LOWERING_OVERLOAD_OP	37	from .lowering_op_list import GENERATE_LIST, GENERATE_LIST2, FALLBACK_LIST, LOWERING_OVERLOAD_OP
38		38
39		39
	@@ -198,6 +198,10 @@ def _register_npu_inductor_fallbacks():
198	def _convert_npu_type(x: TensorBox, dtype: torch.dtype):	198	def _convert_npu_type(x: TensorBox, dtype: torch.dtype):
199	return to_dtype(x, dtype, copy=True)	199	return to_dtype(x, dtype, copy=True)
200		200
		201	+ @register_lowering(_npu_dtype_cast, type_promotion_kind=None)
		202	+ def _convert__npu_type(x: TensorBox, dtype: torch.dtype):
		203	+ return to_dtype(x, dtype, copy=True)
		204	+
201	def var_mean_sum_(x, axis, correction, keepdim, return_mean):	205	def var_mean_sum_(x, axis, correction, keepdim, return_mean):
202	if correction is None:	206	if correction is None:
203	correction = 1	207	correction = 1

Mtorch_npu/_inductor/lowering_fx.py+4-0

	@@ -2223,6 +2223,10 @@ def _register_npu_inductor_fallbacks():
2223	def _convert_npu_type(x: TensorBox, dtype: torch.dtype):	2223	def _convert_npu_type(x: TensorBox, dtype: torch.dtype):
2224	return to_dtype(x, dtype, copy=True)	2224	return to_dtype(x, dtype, copy=True)
2225		2225
		2226	+ @register_lowering(npu._npu_dtype_cast, type_promotion_kind=None)
		2227	+ def _convert__npu_type(x: TensorBox, dtype: torch.dtype):
		2228	+ return to_dtype(x, dtype, copy=True)
		2229	+
2226	def var_mean_sum_(x, axis, correction, keepdim, return_mean):	2230	def var_mean_sum_(x, axis, correction, keepdim, return_mean):
2227	if correction is None:	2231	if correction is None:
2228	correction = 1	2232	correction = 1

Mtorch_npu/_inductor/lowering_op_list.py+2-1

	@@ -1,5 +1,5 @@
1	import torch	1	import torch
2	-from torch_npu import npu_dtype_cast	2	+from torch_npu import npu_dtype_cast, _npu_dtype_cast
3		3
4	aten = torch.ops.aten	4	aten = torch.ops.aten
5	tr_c10d = torch.ops.tr_c10d	5	tr_c10d = torch.ops.tr_c10d
	@@ -56,6 +56,7 @@ GENERATE_LIST = [
56	aten.clamp_max,	56	aten.clamp_max,
57	aten.mean,	57	aten.mean,
58	npu_dtype_cast,	58	npu_dtype_cast,
		59	+ _npu_dtype_cast,
59	aten.select_scatter,	60	aten.select_scatter,
60	aten.slice_scatter,	61	aten.slice_scatter,
61	prims.broadcast_in_dim,	62	prims.broadcast_in_dim,

Mtorch_npu/_inductor/npu_triton_heuristics.py+73-25

	@@ -1,6 +1,7 @@
1	# This file is based on triton_heuristics with heuristics designed for NPU	1	# This file is based on triton_heuristics with heuristics designed for NPU
2	import copy	2	import copy
3	import functools	3	import functools
		4	+from functools import lru_cache
4	import hashlib	5	import hashlib
5	import importlib	6	import importlib
6	import json	7	import json
	@@ -551,6 +552,15 @@ class NPUCachingAutotuner(CachingAutotuner):
551	if self.save_cache_hook:	552	if self.save_cache_hook:
552	self.save_cache_hook(self.launchers[0].config, self.autotune_time_taken_ns)	553	self.save_cache_hook(self.launchers[0].config, self.autotune_time_taken_ns)
553		554
		555	+ @lru_cache(None)
		556	+ def get_fx_graph_dump_path(self):
		557	+ traced_graph_hash = self.inductor_meta.get("traced_graph_hash")
		558	+ dump_dir = self.inductor_meta.get("traced_graph_dir", "")
		559	+ dump_path = os.path.join(dump_dir, traced_graph_hash)
		560	+ if dump_dir == "" or not os.path.exists(dump_path):
		561	+ return None
		562	+ return dump_path
		563	+
554	def get_fx_graph_call(self, auto_fallback=False):	564	def get_fx_graph_call(self, auto_fallback=False):
555	kernel_name = self.inductor_meta.get("kernel_name", "triton_")	565	kernel_name = self.inductor_meta.get("kernel_name", "triton_")
556	traced_graph_hash = self.inductor_meta.get("traced_graph_hash")	566	traced_graph_hash = self.inductor_meta.get("traced_graph_hash")
	@@ -593,8 +603,13 @@ class NPUCachingAutotuner(CachingAutotuner):
593	return fx_graph_call, kernel_name, dump_path, fx_module	603	return fx_graph_call, kernel_name, dump_path, fx_module
594		604
595	def data_dump(self, *args, dump_path=None):	605	def data_dump(self, *args, dump_path=None):
		606	+ dump_path = self.get_fx_graph_dump_path() if dump_path is None else dump_path
		607	+ if dump_path is None:
		608	+ log.warning(f"data dump for kernel {self.get_fn_name()} failed, no valid dump_path is supplied.")
		609	+ return False
596	data_dump_path = os.path.join(dump_path, 'data.pth')	610	data_dump_path = os.path.join(dump_path, 'data.pth')
597	torch.save(args, data_dump_path)	611	torch.save(args, data_dump_path)
		612	+ return True
598		613
599	def get_fn_name(self):	614	def get_fn_name(self):
600	if self.fn_name is not None:	615	if self.fn_name is not None:
	@@ -651,12 +666,11 @@ class NPUCachingAutotuner(CachingAutotuner):
651	return True	666	return True
652		667
653		668
654	- def check_accuracy(self, args, launcher, stream, *kwargs):	669	+ def check_accuracy(self, args, launcher, grid, stream, *kwargs):
655	fx_graph_call, kernel_name, dump_path, fx_module = self.get_fx_graph_call()	670	fx_graph_call, kernel_name, dump_path, fx_module = self.get_fx_graph_call()
656	if not fx_graph_call:	671	if not fx_graph_call:
657	return None	672	return None
658	call_outputs_indices = fx_module.call_args_mapping[fx_module.num_inputs:]	673	call_outputs_indices = fx_module.call_args_mapping[fx_module.num_inputs:]
659	- self.data_dump(*args, dump_path=dump_path)
660		674
661	fx_args = []	675	fx_args = []
662	for idx in fx_module.call_args_mapping:	676	for idx in fx_module.call_args_mapping:
	@@ -668,10 +682,9 @@ class NPUCachingAutotuner(CachingAutotuner):
668		682
669	fx_graph_call(*fx_args)	683	fx_graph_call(*fx_args)
670		684
671	- ret = launcher(	685	+ launcher(
672	*args,	686	*args,
673	**kwargs,	687	**kwargs,
674	- grid=grid,
675	stream=stream,	688	stream=stream,
676	)	689	)
677		690
	@@ -730,16 +743,42 @@ class NPUCachingAutotuner(CachingAutotuner):
730	torch.save(dump_args, f"{dump_path}/{idx}_{fn_name}_after.pt")	743	torch.save(dump_args, f"{dump_path}/{idx}_{fn_name}_after.pt")
731	return result	744	return result
732		745
		746	+ def maybe_run_debug(self, args, grid_, stream, launcher, *kwargs):
		747	+ kernel_name = self.get_fn_name()
		748	+ log.info(f"Try to run debug mode for kernel {kernel_name}.")
		749	+ if npu_config.dump_fx_graph:
		750	+ _ = self.data_dump(*args)
		751	+
		752	+ if npu_config.check_accuracy:
		753	+ if self.check_accuracy(args, launcher=launcher, grid=grid_, stream=stream, *kwargs):
		754	+ return "check_accuracy"
		755	+ elif npu_config.force_fallback_kernel_id:
		756	+ fallback_result = self.fallback_to_fx(args, launcher=launcher, grid_=grid_, stream=stream, *kwargs)
		757	+ if fallback_result is not None:
		758	+ log.debug(f"fallback kernel {self.get_fn_name()} to fx graph call.")
		759	+ return "force_fallback_kernel_id"
		760	+ else:
		761	+ log.warning(f"kernel {self.get_fn_name()} could not fallback to fx.")
		762	+ elif npu_config.aot_inductor.debug_kernel_in_run:
		763	+ _ = self.debug_kernel_in_run(args, launcher=launcher, grid_=grid_, stream=stream, *kwargs)
		764	+ return "debug_kernel_in_run"
		765	+
		766	+ log.info(f"No debug mode is activated for kernel {kernel_name}.")
		767	+ return None
733		768
734	def run(	769	def run(
735	self, args, stream, benchmark_run=False, *kwargs	770	self, args, stream, benchmark_run=False, *kwargs
736	): # type:ignore[override]	771	): # type:ignore[override]
737	if self.triton_interpret:	772	if self.triton_interpret:
738	args, grid = self._interpret_args_grid(args, self.configs[0])	773	args, grid = self._interpret_args_grid(args, self.configs[0])
		774	+ copied_kwargs = copy.copy(self.configs[0].kwargs)
		775	+ copied_kwargs.pop('split_axis', None)
		776	+ copied_kwargs.pop('split_blocks', None)
		777	+
739	return self.fn[grid](	778	return self.fn[grid](
740	*args,	779	*args,
741	**kwargs,	780	**kwargs,
742	- **self.configs[0].kwargs,	781	+ **copied_kwargs,
743	)	782	)
744		783
745	if hasattr(self.launchers[0], "fallback"):	784	if hasattr(self.launchers[0], "fallback"):
	@@ -772,26 +811,11 @@ class NPUCachingAutotuner(CachingAutotuner):
772	if self.dump_launch_params:	811	if self.dump_launch_params:
773	_dump_launch_params(args, kwargs, launcher, self.fn.__name__)	812	_dump_launch_params(args, kwargs, launcher, self.fn.__name__)
774		813
775	- if npu_config.check_accuracy:	814	+ _, grid = self._interpret_args_grid(args, launcher.config)
776	- if self.check_accuracy(args, launcher=launcher, stream=stream, *kwargs):	815	+ debug_mode = self.maybe_run_debug(args, grid_=grid, stream=stream, launcher=launcher, *kwargs)
777	- return	816	+ if debug_mode:
778	-	817	+ log.info(f"Kernel {self.get_fn_name()} goes into {debug_mode} and return.")
779	- elif npu_config.dump_fx_graph:	818	+ return
780	- fx_graph_call, kernel_name, dump_path, _ = self.get_fx_graph_call()
781	- if not fx_graph_call:
782	- log.warning(f"data dump for kernel {kernel_name} failed!")
783	- else:
784	- self.data_dump(*args, dump_path=dump_path)
785	-
786	- elif npu_config.force_fallback_kernel_id:
787	- fallback_result = self.fallback_to_fx(args, launcher=launcher, stream=stream, *kwargs)
788	- if fallback_result is not None:
789	- log.debug(f"fallback kernel {self.get_fn_name()} to fx graph call.")
790	- return
791	- else:
792	- log.warning(f"kernel {self.get_fn_name()} could not fallback to fx.")
793	- elif npu_config.aot_inductor.debug_kernel_in_run:
794	- return self.debug_kernel_in_run(args, launcher=launcher, stream=stream, *kwargs)
795		819
796	# it is faster than entering and exiting a context manager, even if the context	820	# it is faster than entering and exiting a context manager, even if the context
797	# manager is a nullcontext.	821	# manager is a nullcontext.
	@@ -818,6 +842,30 @@ class NPUCachingAutotuner(CachingAutotuner):
818	stream=stream,	842	stream=stream,
819	)	843	)
820		844
		845	+ def _interpret_args_grid(
		846	+ self, args: tuple[Any, ...], cfg: Config
		847	+ ) -> tuple[tuple[Any, ...], tuple[int, int, int]]:
		848	+
		849	+ numels = [
		850	+ arg
		851	+ for arg in self.fn.arg_names
		852	+ if "_numel" in arg
		853	+ ]
		854	+ grid = GridExprNpu.from_meta_and_set_numel(self.inductor_meta, cfg, numels).eval_slow(
		855	+ dict(
		856	+ zip(
		857	+ [
		858	+ *self.fn.arg_names,
		859	+ *self.inductor_meta.get("extra_launcher_args", ()),
		860	+ ],
		861	+ args,
		862	+ )
		863	+ )
		864	+ )
		865	+ if self.inductor_meta.get("extra_launcher_args"):
		866	+ args = args[: -len(self.inductor_meta["extra_launcher_args"])]
		867	+ return args, grid
		868	+
821		869
822	class NPUDebugAutotuner(NPUCachingAutotuner):	870	class NPUDebugAutotuner(NPUCachingAutotuner):
823	def __init__(self, args, regex_filter="", *kwargs):	871	def __init__(self, args, regex_filter="", *kwargs):

Mtorch_npu/_inductor/utils.py+3-1

	@@ -73,4 +73,6 @@ def patch_has_triton():
73	return is_device_compatible_with_triton()	73	return is_device_compatible_with_triton()
74		74
75	torch.utils._triton.has_triton = has_triton	75	torch.utils._triton.has_triton = has_triton
76	- torch._inductor.scheduler.has_triton = has_triton	76	+ torch._inductor.scheduler.has_triton = has_triton
		77	+
		78	+

Mtorch_npu/utils/_dynamo.py+59-2

@@ @@ -1,4 +1,5 @@ @@
 import inspect
++import sys
 from typing import Dict, List
 import torch
 class NPUTorchCtxManagerClassVariable(TorchCtxManagerClassVariable):
     def call_function(self, tx, args, kwargs):
++        return NPUAutocastModeVariable.create(self.value, args, kwargs)
 class NPUAutocastModeVariable(AutocastModeVariable):
         return TensorVariable.call_method_raw(self, tx, name, args, kwargs)
++class _InductorNpuRegistry:
++    _disabled_register = False
++    _has_inited = False
++
++    @classmethod
++    def register_inductor_npu(cls):
++        if cls.has_initialized() or cls._disabled_register:
++            return
++        from torch_npu import _inductor
++        cls._has_inited = True
++
++    @classmethod
++    def disable_register(cls):
++        cls._disabled_register = True
++
++    @classmethod
++    def enable_register(cls):
++        cls._disabled_register = False
++
++    @classmethod
++    def has_initialized(cls):
++        if cls._has_inited:
++            return True
++        # Maybe initialized by call `import torch_npu._inductor` manually.
++        if 'torch_npu._inductor' in sys.modules:
++            cls._has_inited = True
++        return cls._has_inited
++
++
++def is_inductor_npu_initialized():
++    return _InductorNpuRegistry.has_initialized()
++
++
++def disable_register_inductor_npu():
++    _InductorNpuRegistry.disable_register()
++
++
++def enable_register_inductor_npu():
++    _InductorNpuRegistry.enable_register()
++
++
++def register_inductor_npu():
++    _InductorNpuRegistry.register_inductor_npu()
++
++
++def patch_inductor_wrapper():
++    from torch import _TorchCompileInductorWrapper
++    src_call = _TorchCompileInductorWrapper.__call__
++
++    def new_call(self, model_, inputs_):
++        register_inductor_npu()
++        return src_call(self, model_, inputs_)
++
++    _TorchCompileInductorWrapper.__call__ = new_call
++
++
 def patch_dynamo_optimize():
     src_optimize = optimize
     TensorVariable.call_method_raw = TensorVariable.call_method
     TensorVariable.call_method = TensorVariable_call_method
     patch_dynamo_optimize()
++    patch_inductor_wrapper()