同步f2209ee61066ae89d4efd1782a30b6235e0e0548-pytorch-AtomGit

Merged

同步f2209ee61066ae89d4efd1782a30b6235e0e0548 #4

AtomGit-BotAdd on Jul 5, 2021Merged frommaster into master

同步f2209ee61066ae89d4efd1782a30b6235e0e0548 #4

Merged

AtomGit-BotAdd on Jul 5, 2021

Merged frommasterintomaster

Collapse

patch
- npu.patch
src
- aten / src
  - ATen / native
    - native_functions.yaml
    - npu
      - AcosKernelNpu.cpp
      - AddbmmKernelNpu.cpp
      - AddcdivKernelNpu.cpp
      - AddcmulKernelNpu.cpp
      - AddmvKernelNpu.cpp
      - AddrKernelNpu.cpp
      - AffineGridGeneratorBackwardKernelNpu.cpp
      - AffineGridGeneratorKernelNpu.cpp
      - AllKernelNpu.cpp
      - AnyKernelNpu.cpp
      - ArangeKernelNpu.cpp
      - ArgminKernelNpu.cpp
      - ArgsortKernelNpu.cpp
      - AsinKernelNpu.cpp
      - Atan2KernelNpu.cpp
      - AtanKernelNpu.cpp
      - BatchNMSKernelNpu.cpp
      - BernoulliKernelNpu.cpp
      - BinaryCrossEntropyBackwardKernelNpu.cpp
      - BinaryCrossEntropyKernelNpu.cpp
      - BinaryCrossEntropyWithLogitsBackwardKernelNpu.cpp
      - BinaryCrossEntropyWithLogitsKernelNpu.cpp
      - BitwiseAndKernelNpu.cpp
      - BitwiseNotKernelNpu.cpp
      - BitwiseOrKernelNpu.cpp
      - BitwiseXorKernelNpu.cpp
      - BmmKernelNpu.cpp
      - BmmV2KernelNpu.cpp
      - BoundingBoxDecodeKernelNpu.cpp
      - BoundingBoxEncodeKernelNpu.cpp
      - BroadcastKernelNpu.cpp
      - CastKernelNpu.cpp
      - CatKernelNpu.cpp
      - CdistBackwardKernelNpu.cpp
      - CeilKernelNpu.cpp
      - CeluKernelNpu.cpp
      - CholeskyKernelNpu.cpp
      - ClampKernelNpu.cpp
      - ConfusionTransposeKernelNpu.cpp
      - ConvTbcKernelNpu.cpp
      - CosKernelNpu.cpp
      - CoshKernelNpu.cpp
      - CrossKernelNpu.cpp
      - CtcLossBackwardKernelNpu.cpp
      - CtcLossKernelNpu.cpp
      - CumprodKernelNpu.cpp
      - DiagKernelNpu.cpp
      - DivKernelNpu.cpp
      - DotKernelNpu.cpp
      - DropoutBackwardKernelNpu.cpp
      - DropoutV2BackwardKernelNpu.cpp
      - DropoutV2KernelNpu.cpp
      - EmbeddingKernelNpu.cpp
      - EqKernelNpu.cpp
      - ExpKernelNpu.cpp
      - FillKernelNpu.cpp
      - FlipKernelNpu.cpp
      - FloorDivideKernelNpu.cpp
      - FloorKernelNpu.cpp
      - FmodKernelNpu.cpp
      - GatherKernelNpu.cpp
      - GeKernelNpu.cpp
      - GeluKernelNpu.cpp
      - GerKernelNpu.cpp
      - GridSampler3dBackwardKernelNpu.cpp
      - GridSampler3dKernelNpu.cpp
      - GtKernelNpu.cpp
      - HardsigmoidKernelNpu.cpp
      - HardtanhBackwardKernelNpu.cpp
      - HardtanhKernelNpu.cpp
      - IfmrKernelNpu.cpp
      - Im2colBackwardKernelNpu.cpp
      - Im2colKernelNpu.cpp
      - IndexAddKernelNpu.cpp
      - IndexFillDKernelNpu.cpp
      - IndexPutKernelNpu.cpp
      - IndexSelectKernelNpu.cpp
      - IndexingKernelNpu.cpp
      - InstanceNormKernelNpu.cpp
      - InverseKernelNpu.cpp
      - L1LossBackwardKernelNpu.cpp
      - L1lossKernelNpu.cpp
      - LeKernelNpu.cpp
      - LeakyReluBackwardKernelNpu.cpp
      - LeakyReluKernelNpu.cpp
      - LinspaceKernelNpu.cpp
      - Log10KernelNpu.cpp
      - Log1pKernelNpu.cpp
      - Log2KernelNpu.cpp
      - LogKernelNpu.cpp
      - LogSigmoidKernelNpu.cpp
      - LogSoftmaxBackwardKernelNpu.cpp
      - LogSoftmaxKernelNpu.cpp
      - LogSpaceKernelNpu.cpp
      - LogicalNotKernelNpu.cpp
      - LstmKernelNpu.cpp
      - LtKernelNpu.cpp
      - MaskedFillKernelNpu.cpp
      - MaskedFillRangeKernelNpu.cpp
      - MaxKernelNpu.cpp
      - MaxV1KernelNpu.cpp
      - MeanKernelNpu.cpp
      - MedianKernelNpu.cpp
      - MinKernelNpu.cpp
      - MseLossBackwardKernelNpu.cpp
      - MseLossKernelNpu.cpp
      - MulKernelNpu.cpp
      - MvKernelNpu.cpp
      - NeKernelNpu.cpp
      - NonzeroKernelNpu.cpp
      - NormKernelNpu.cpp
      - NormalKernelNpu.cpp
      - OneHotKernelNpu.cpp
      - OnesLikeKernelNpu.cpp
      - PadKernelNpu.cpp
      - PowKernelNpu.cpp
      - PreluBackwardKernelNpu.cpp
      - PreluKernelNpu.cpp
      - ProdKernelNpu.cpp
      - PtIouKernelNpu.cpp
      - QrKernelNpu.cpp
      - RandomChoiceWithMaskKernelNpu.cpp
      - RandomKernelNpu.cpp
      - RangeKernelNpu.cpp
      - ReciprocalKernelNpu.cpp
      - RemainderKernelNpu.cpp
      - RepeatKernelNpu.cpp
      - RollKernelNpu.cpp
      - RoundKernelNpu.cpp
      - RsqrtKernelNpu.cpp
      - RsubKernelNpu.cpp
      - ScatterAddKernelNpu.cpp
      - ScatterKernelNpu.cpp
      - ScatterV1KernelNpu.cpp
      - SeluKernelNpu.cpp
      - SigmoidBackwardKernelNpu.cpp
      - SigmoidKernelNpu.cpp
      - SignKernelNpu.cpp
      - SinhKernelNpu.cpp
      - SliceKernelNpu.cpp
      - SmoothL1LossKernelNpu.cpp
      - SoftMarginLossBackwardKernelNpu.cpp
      - SoftmaxBackwardKernelNpu.cpp
      - SoftmaxKernelNpu.cpp
      - SortKernelNpu.cpp
      - SortWithoutIndicesKernelNpu.cpp
      - StackKernelNpu.cpp
      - StdKernelNpu.cpp
      - SubKernelNpu.cpp
      - SubSampleKernelNpu.cpp
      - SumKernelNpu.cpp
      - TanKernelNpu.cpp
      - TanhBackwardKernelNpu.cpp
      - TanhKernelNpu.cpp
      - ThresholdBackwardKernelNpu.cpp
      - ThresholdKernelNpu.cpp
      - TopKKernelNpu.cpp
      - TransposeKernelNpu.cpp
      - TrilKernelNpu.cpp
      - TruncKernelNpu.cpp
      - UniformKernelNpu.cpp
      - UpSampleNearest3dKernelNpu.cpp
      - UpsampleBilinear2dKernelNpu.cpp
      - UpsampleNearest1dBackwardKernelNpu.cpp
      - UpsampleNearest1dKernelNpu.cpp
      - UpsampleNearest2dBackwardKernelNpu.cpp
      - UpsampleNearest2dKernelNpu.cpp
      - UpsampleNearest3dBackwardKernelNpu.cpp
      - UpsampleTrilinear3dBackwardKernelNpu.cpp
      - UpsampleTrilinear3dKernelNpu.cpp
      - WhereKernelNpu.cpp
      - ZerosLikeKernelNpu.cpp
      - _Unique2KernelNpu.cpp
      - __And__KernelNpu.cpp
      - __Ior__KernelNpu.cpp
      - __Or__KernelNpu.cpp
      - common
        MatmulByBmmV2KernelNpu.cpp
      - convolution
        Conv3dKernelNpu.cpp
      - frame
        OpCommandBase.h
        OpDynamicParamMaker.h
        OpParamMaker.cpp
        OpParamMaker.h
      - interface
        EnvVariables.cpp
        EnvVariables.h
      - loss
        MultilabelMarginLossKernelNpu.cpp
        NLLLoss2dKernelNpu.cpp
        NLLLossBackwardKernelNpu.cpp
        NLLLossKernelNpu.cpp
      - pooling
        AdaptiveAvgPool2dBackwardKernelNpu.cpp
        AdaptiveAvgPool3dBackwardKernelNpu.cpp
        AvgPool2dKernelNpu.cpp
        MaxPool3dWithIndicesBackwardKernelNpu.cpp
        MaxPool3dWithIndicesKernelNpu.cpp
      - utils
        CalcuOpUtil.cpp
        KernelNpuOutputSize.cpp
        KernelNpuOutputSize.h
        NpuProfilingDispatch.cpp
        NpuProfilingDispatch.h
        NpuUtils.cpp
        NpuUtils.h
        OpPipeWithMultiOut.h
        OpPreparation.cpp
        OpPreparation.h
        OpTemplate.cpp
  - THNPU
    - THNPUCachingHostAllocator.cpp
- c10 / npu
  - NPUCachingAllocator.cpp
  - NPUEventManager.cpp
  - NPUStream.cpp
  - OptionsManager.h
  - interface
    - AclInterface.cpp
    - AclInterface.h
  - register
    - FunctionLoader.cpp
- third_party / acl
  - inc / acl
    - acl.h
    - acl_base.h
  - libs
    - acl.cpp
- tools / autograd
  - derivatives.yaml
- torch
  - csrc / autograd
    - profiler_npu.cpp
  - npu
    - __init__.py
    - npu_frontend_enhance.py
test / test_npu
- test_conv_tbc_backward.py
- test_instance_norm.py
- test_instancenorm.py
- test_multilabel_margin_loss_backward.py
- test_network_ops
  - run_tests.py
  - test__Ixor__.py
  - test___iand__.py
  - test__nnpack_spatial_convolution.py
  - test_acos.py
  - test_adaptive_avg_pool2d_backward.py
  - test_adaptive_max_pool2d_backward.py
  - test_addbmm.py
  - test_addcdiv.py
  - test_addmv.py
  - test_addr.py
  - test_affine_grid_generator_backward.py
  - test_asin.py
  - test_bartlett_window.py
  - test_batch_norm.py
  - test_binary_cross_entropy_with_logits.py
  - test_bitwise_not.py
  - test_blackman_window.py
  - test_bmm_v2.py
  - test_cast_Byte.py
  - test_cast_Char.py
  - test_cast_Float.py
  - test_cast_Half.py
  - test_cast_Int.py
  - test_cast_Long.py
  - test_cast_Short.py
  - test_cdist.py
  - test_cdist_backward.py
  - test_celu.py
  - test_conv_tbc.py
  - test_conv_transpose2d.py
  - test_convolution_backward_weight.py
  - test_convolution_transpose_backward_weight.py
  - test_copy.py
  - test_cos.py
  - test_cosh.py
  - test_cosinesimilarity.py
  - test_cross.py
  - test_cudnn_convolution_backward_bias.py
  - test_cudnn_convolution_transpose_backward_bias.py
  - test_cumprod.py
  - test_cumsum.py
  - test_dim_arange.py
  - test_diml.py
  - test_dirichlet_grad.py
  - test_dot.py
  - test_elu.py
  - test_embedding.py
  - test_equal.py
  - test_erf.py
  - test_erfinv.py
  - test_expm1.py
  - test_eye.py
  - test_floordivide.py
  - test_frac.py
  - test_frobenius_norm.py
  - test_full_like.py
  - test_gelu.py
  - test_glu.py
  - test_glugrad.py
  - test_grid_sampler_2d.py
  - test_grid_sampler_2d_backward.py
  - test_group_norm.py
  - test_hamming_window.py
  - test_hammingwindow.py
  - test_hanning_window.py
  - test_hard_sigmoid_backward.py
  - test_hardshrink.py
  - test_hardshrink_backward.py
  - test_hardsigmoid.py
  - test_hinge_embedding_loss.py
  - test_im2col.py
  - test_index_fill_d.py
  - test_index_select.py
  - test_isclose.py
  - test_kthvalue.py
  - test_l1_loss.py
  - test_l1_loss_backward.py
  - test_leaky_relu.py
  - test_log10.py
  - test_log1p.py
  - test_log_sigmoid.py
  - test_log_sigmoid_backward.py
  - test_logical_not.py
  - test_logsigmoid.py
  - test_logsigmoidforward.py
  - test_logsumexp.py
  - test_masked_fill_range.py
  - test_max_pool2d_backward.py
  - test_miopen_batch_norm.py
  - test_miopen_batch_norm_backward.py
  - test_miopen_convolution.py
  - test_miopen_convolution_backward.py
  - test_miopen_convolution_backward_bias.py
  - test_miopen_convolution_backward_input.py
  - test_miopen_convolution_backward_weight.py
  - test_mkldnn_adaptive_avg_pool2d.py
  - test_mkldnn_convolution_backward.py
  - test_mkldnn_convolution_backward_input.py
  - test_mkldnn_convolution_backward_weights.py
  - test_multilabel_margin_loss.py
  - test_multinomial.py
  - test_narrow_copy.py
  - test_nllloss2d.py
  - test_norm_except_dim.py
  - test_norm_ext.py
  - test_one_hot.py
  - test_ones.py
  - test_pdist.py
  - test_pixel_shuffle.py
  - test_pooling.py
  - test_prelu.py
  - test_prelu_backward.py
  - test_quantize_per_channel.py
  - test_quantize_per_tensor.py
  - test_real.py
  - test_renorm.py
  - test_repeat_interleave.py
  - test_resize_.py
  - test_roll.py
  - test_selu.py
  - test_sinh.py
  - test_slow_conv_dilated2d.py
  - test_slow_conv_dilated2d_backward.py
  - test_slow_conv_transpose2d.py
  - test_slow_conv_transpose2d_backward.py
  - test_slow_conv_transpose3d.py
  - test_soft_margin_loss.py
  - test_softmax_backward.py
  - test_softshrink.py
  - test_softshrink_backward.py
  - test_split_with_sizes.py
  - test_square.py
  - test_std.py
  - test_sum_to_size.py
  - test_take.py
  - test_tan.py
  - test_tensor_npu.py
  - test_thnn_conv_depthwise2d_backward.py
  - test_thnn_conv_depthwise2d_forward.py
  - test_threshold_grad_v2_d.py
  - test_threshold_v2_d.py
  - test_trapz_dx.py
  - test_trapz_x.py
  - test_true_divide.py
  - test_trunc.py
  - test_unbind_copy_contiguous.py
  - test_upsample_bilinear2d.py
  - test_upsample_linear1d.py
  - test_upsample_nearest2d.py
  - test_upsample_nearest2d_backward.py
  - test_upsample_nearest3d.py
  - test_upsample_nearest3d_backward.py
  - test_upsample_trilinear3d.py
  - test_upsample_trilinear3d_backward.py
- test_trace.py

Total: 381 files changed+3298-3098

Mpatch/npu.patch+629-585

Msrc/aten/src/ATen/native/native_functions.yaml+65-46

	@@ -606,7 +606,7 @@
606	dispatch:	606	dispatch:
607	CPU: bernoulli_scalar_cpu_	607	CPU: bernoulli_scalar_cpu_
608	CUDA: bernoulli_scalar_cuda_	608	CUDA: bernoulli_scalar_cuda_
609	- supports_named_tensor: True	609	+ supports_named_tensor: True
610	npu_dispatch:	610	npu_dispatch:
611	NPU: bernoulli_npu_	611	NPU: bernoulli_npu_
612		612
	@@ -1040,7 +1040,7 @@
1040	CUDA: _cosh_out_cuda	1040	CUDA: _cosh_out_cuda
1041	npu_dispatch:	1041	npu_dispatch:
1042	NPU: cosh_out_npu	1042	NPU: cosh_out_npu
1043	-	1043	+
1044	- func: cosine_embedding_loss(Tensor input1, Tensor input2, Tensor target, float margin=0.0, int reduction=Mean) -> Tensor	1044	- func: cosine_embedding_loss(Tensor input1, Tensor input2, Tensor target, float margin=0.0, int reduction=Mean) -> Tensor
1045	use_c10_dispatcher: full	1045	use_c10_dispatcher: full
1046		1046
	@@ -1271,7 +1271,7 @@
1271	- func: det(Tensor self) -> Tensor	1271	- func: det(Tensor self) -> Tensor
1272	use_c10_dispatcher: full	1272	use_c10_dispatcher: full
1273	variants: function, method	1273	variants: function, method
1274	-	1274	+
1275	- func: diag_embed(Tensor self, int offset=0, int dim1=-2, int dim2=-1) -> Tensor	1275	- func: diag_embed(Tensor self, int offset=0, int dim1=-2, int dim2=-1) -> Tensor
1276	use_c10_dispatcher: full	1276	use_c10_dispatcher: full
1277	variants: function, method	1277	variants: function, method
	@@ -1381,7 +1381,7 @@
1381	CUDA: embedding_renorm_cuda_	1381	CUDA: embedding_renorm_cuda_
1382	npu_dispatch:	1382	npu_dispatch:
1383	NPU: embedding_renorm_npu_	1383	NPU: embedding_renorm_npu_
1384	-	1384	+
1385	- func: embedding_sparse_backward(Tensor grad, Tensor indices, int num_weights, int padding_idx, bool scale_grad_by_freq) -> Tensor	1385	- func: embedding_sparse_backward(Tensor grad, Tensor indices, int num_weights, int padding_idx, bool scale_grad_by_freq) -> Tensor
1386	use_c10_dispatcher: full	1386	use_c10_dispatcher: full
1387		1387
	@@ -1657,7 +1657,7 @@
1657	SparseCPU: floor_divide_sparse	1657	SparseCPU: floor_divide_sparse
1658	SparseCUDA: floor_divide_sparse	1658	SparseCUDA: floor_divide_sparse
1659	supports_named_tensor: True	1659	supports_named_tensor: True
1660	- npu_dispatch:	1660	+ npu_dispatch:
1661	NPU: floor_divide_npu	1661	NPU: floor_divide_npu
1662		1662
1663	- func: floor_divide_.Tensor(Tensor(a!) self, Tensor other) -> Tensor(a!)	1663	- func: floor_divide_.Tensor(Tensor(a!) self, Tensor other) -> Tensor(a!)
	@@ -1668,7 +1668,7 @@
1668	SparseCPU: floor_divide_sparse_	1668	SparseCPU: floor_divide_sparse_
1669	SparseCUDA: floor_divide_sparse_	1669	SparseCUDA: floor_divide_sparse_
1670	supports_named_tensor: True	1670	supports_named_tensor: True
1671	- npu_dispatch:	1671	+ npu_dispatch:
1672	NPU: floor_divide_npu_	1672	NPU: floor_divide_npu_
1673		1673
1674	- func: floor_divide.out(Tensor self, Tensor other, *, Tensor(a!) out) -> Tensor(a!)	1674	- func: floor_divide.out(Tensor self, Tensor other, *, Tensor(a!) out) -> Tensor(a!)
	@@ -1678,13 +1678,13 @@
1678	SparseCPU: floor_divide_out_sparse_zerodim	1678	SparseCPU: floor_divide_out_sparse_zerodim
1679	SparseCUDA: floor_divide_out_sparse_zerodim	1679	SparseCUDA: floor_divide_out_sparse_zerodim
1680	supports_named_tensor: True	1680	supports_named_tensor: True
1681	- npu_dispatch:	1681	+ npu_dispatch:
1682	NPU: floor_divide_out_npu	1682	NPU: floor_divide_out_npu
1683		1683
1684	- func: floor_divide.Scalar(Tensor self, Scalar other) -> Tensor	1684	- func: floor_divide.Scalar(Tensor self, Scalar other) -> Tensor
1685	variants: function, method	1685	variants: function, method
1686	supports_named_tensor: True	1686	supports_named_tensor: True
1687	- npu_dispatch:	1687	+ npu_dispatch:
1688	NPU: floor_divide_npu	1688	NPU: floor_divide_npu
1689		1689
1690	- func: floor_divide_.Scalar(Tensor(a!) self, Scalar other) -> Tensor(a!)	1690	- func: floor_divide_.Scalar(Tensor(a!) self, Scalar other) -> Tensor(a!)
	@@ -1793,15 +1793,15 @@
1793	- func: hamming_window.periodic(int window_length, bool periodic, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor	1793	- func: hamming_window.periodic(int window_length, bool periodic, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor
1794	npu_dispatch:	1794	npu_dispatch:
1795	NPU: hamming_window_npu	1795	NPU: hamming_window_npu
1796	-	1796	+
1797	- func: hamming_window.periodic_alpha(int window_length, bool periodic, float alpha, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor	1797	- func: hamming_window.periodic_alpha(int window_length, bool periodic, float alpha, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor
1798	npu_dispatch:	1798	npu_dispatch:
1799	NPU: hamming_window_npu	1799	NPU: hamming_window_npu
1800	-	1800	+
1801	- func: hamming_window.periodic_alpha_beta(int window_length, bool periodic, float alpha, float beta, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor	1801	- func: hamming_window.periodic_alpha_beta(int window_length, bool periodic, float alpha, float beta, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor
1802	npu_dispatch:	1802	npu_dispatch:
1803	NPU: hamming_window_npu	1803	NPU: hamming_window_npu
1804	-	1804	+
1805	- func: hinge_embedding_loss(Tensor self, Tensor target, float margin=1.0, int reduction=Mean) -> Tensor	1805	- func: hinge_embedding_loss(Tensor self, Tensor target, float margin=1.0, int reduction=Mean) -> Tensor
1806	use_c10_dispatcher: full	1806	use_c10_dispatcher: full
1807		1807
	@@ -1995,7 +1995,7 @@
1995	dispatch:	1995	dispatch:
1996	CPU: kthvalue_out_cpu	1996	CPU: kthvalue_out_cpu
1997	CUDA: kthvalue_out_cuda	1997	CUDA: kthvalue_out_cuda
1998	- npu_dispatch:	1998	+ npu_dispatch:
1999	NPU: kthvalue_out_npu	1999	NPU: kthvalue_out_npu
2000		2000
2001	- func: kthvalue.dimname(Tensor self, int k, Dimname dim, bool keepdim=False) -> (Tensor values, Tensor indices)	2001	- func: kthvalue.dimname(Tensor self, int k, Dimname dim, bool keepdim=False) -> (Tensor values, Tensor indices)
	@@ -2006,7 +2006,7 @@
2006		2006
2007	- func: kthvalue.dimname_out(Tensor self, int k, Dimname dim, bool keepdim=False, *, Tensor(a!) values, Tensor(b!) indices) -> (Tensor(a!) values, Tensor(b!) indices)	2007	- func: kthvalue.dimname_out(Tensor self, int k, Dimname dim, bool keepdim=False, *, Tensor(a!) values, Tensor(b!) indices) -> (Tensor(a!) values, Tensor(b!) indices)
2008	supports_named_tensor: True	2008	supports_named_tensor: True
2009	- npu_dispatch:	2009	+ npu_dispatch:
2010	NPU: kthvalue_out_npu	2010	NPU: kthvalue_out_npu
2011		2011
2012	- func: layer_norm(Tensor input, int[] normalized_shape, Tensor? weight=None, Tensor? bias=None, float eps=1e-05, bool cudnn_enable=True) -> Tensor	2012	- func: layer_norm(Tensor input, int[] normalized_shape, Tensor? weight=None, Tensor? bias=None, float eps=1e-05, bool cudnn_enable=True) -> Tensor
	@@ -2090,16 +2090,22 @@
2090	use_c10_dispatcher: full	2090	use_c10_dispatcher: full
2091	supports_named_tensor: True	2091	supports_named_tensor: True
2092	variants: function, method	2092	variants: function, method
		2093	+ npu_dispatch:
		2094	+ NPU: log10_npu
2093		2095
2094	- func: log10_(Tensor(a!) self) -> Tensor(a!)	2096	- func: log10_(Tensor(a!) self) -> Tensor(a!)
2095	supports_named_tensor: True	2097	supports_named_tensor: True
2096	variants: function, method	2098	variants: function, method
		2099	+ npu_dispatch:
		2100	+ NPU: log10_npu_
2097		2101
2098	- func: log10.out(Tensor self, *, Tensor(a!) out) -> Tensor(a!)	2102	- func: log10.out(Tensor self, *, Tensor(a!) out) -> Tensor(a!)
2099	supports_named_tensor: True	2103	supports_named_tensor: True
2100	dispatch:	2104	dispatch:
2101	CPU: log10_out	2105	CPU: log10_out
2102	CUDA: log10_out	2106	CUDA: log10_out
		2107	+ npu_dispatch:
		2108	+ NPU: log10_out_npu
2103		2109
2104	- func: log1p(Tensor self) -> Tensor	2110	- func: log1p(Tensor self) -> Tensor
2105	use_c10_dispatcher: full	2111	use_c10_dispatcher: full
	@@ -2730,14 +2736,14 @@
2730	use_c10_dispatcher: full	2736	use_c10_dispatcher: full
2731	npu_dispatch:	2737	npu_dispatch:
2732	NPU: _pdist_forward_npu	2738	NPU: _pdist_forward_npu
2733	-	2739	+
2734	- func: _pdist_backward(Tensor grad, Tensor self, float p, Tensor pdist) -> Tensor	2740	- func: _pdist_backward(Tensor grad, Tensor self, float p, Tensor pdist) -> Tensor
2735	use_c10_dispatcher: full	2741	use_c10_dispatcher: full
2736		2742
2737	- func: cosine_similarity(Tensor input, Tensor input2, int dim=1, float eps=1e-08) -> Tensor	2743	- func: cosine_similarity(Tensor input, Tensor input2, int dim=1, float eps=1e-08) -> Tensor
2738	use_c10_dispatcher: full	2744	use_c10_dispatcher: full
2739	variants: function	2745	variants: function
2740	-	2746	+
2741	- func: permute(Tensor(a) self, int[] dims) -> Tensor(a)	2747	- func: permute(Tensor(a) self, int[] dims) -> Tensor(a)
2742	variants: method # This is method-only to match the previous tensor API. In the future we could make this a function too.	2748	variants: method # This is method-only to match the previous tensor API. In the future we could make this a function too.
2743		2749
	@@ -2753,7 +2759,7 @@
2753		2759
2754	- func: pixel_shuffle(Tensor self, int upscale_factor) -> Tensor	2760	- func: pixel_shuffle(Tensor self, int upscale_factor) -> Tensor
2755	use_c10_dispatcher: full	2761	use_c10_dispatcher: full
2756	-	2762	+
2757	- func: is_pinned(Tensor self) -> bool	2763	- func: is_pinned(Tensor self) -> bool
2758	use_c10_dispatcher: full	2764	use_c10_dispatcher: full
2759	variants: method	2765	variants: method
	@@ -2770,7 +2776,7 @@
2770	- func: poisson_nll_loss(Tensor input, Tensor target, bool log_input, bool full, float eps, int reduction) -> Tensor	2776	- func: poisson_nll_loss(Tensor input, Tensor target, bool log_input, bool full, float eps, int reduction) -> Tensor
2771	use_c10_dispatcher: full	2777	use_c10_dispatcher: full
2772	variants: function	2778	variants: function
2773	-	2779	+
2774	- func: scalar_tensor(Scalar s, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor	2780	- func: scalar_tensor(Scalar s, *, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor
2775		2781
2776	- func: rand.names(int[] size, *, Dimname[]? names, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor	2782	- func: rand.names(int[] size, *, Dimname[]? names, ScalarType? dtype=None, Layout? layout=None, Device? device=None, bool? pin_memory=None) -> Tensor
	@@ -3011,7 +3017,7 @@
3011	CUDA: gelu_cuda	3017	CUDA: gelu_cuda
3012	npu_dispatch:	3018	npu_dispatch:
3013	NPU: gelu_npu	3019	NPU: gelu_npu
3014	-	3020	+
3015	- func: gelu_backward(Tensor grad, Tensor self) -> Tensor	3021	- func: gelu_backward(Tensor grad, Tensor self) -> Tensor
3016	use_c10_dispatcher: full	3022	use_c10_dispatcher: full
3017	python_module: nn	3023	python_module: nn
	@@ -3075,8 +3081,12 @@
3075		3081
3076	- func: celu(Tensor self, Scalar alpha=1.0) -> Tensor	3082	- func: celu(Tensor self, Scalar alpha=1.0) -> Tensor
3077	use_c10_dispatcher: full	3083	use_c10_dispatcher: full
		3084	+ npu_dispatch:
		3085	+ NPU: celu_npu
3078		3086
3079	- func: celu_(Tensor(a!) self, Scalar alpha=1.0) -> Tensor(a!)	3087	- func: celu_(Tensor(a!) self, Scalar alpha=1.0) -> Tensor(a!)
		3088	+ npu_dispatch:
		3089	+ NPU: celu_npu_
3080		3090
3081	- func: sigmoid(Tensor self) -> Tensor	3091	- func: sigmoid(Tensor self) -> Tensor
3082	use_c10_dispatcher: full	3092	use_c10_dispatcher: full
	@@ -3423,14 +3433,14 @@
3423	npu_dispatch:	3433	npu_dispatch:
3424	NPU: prod_out_npu	3434	NPU: prod_out_npu
3425	#NPU: prod_out_npu_ext	3435	#NPU: prod_out_npu_ext
3426	-	3436	+
3427	- func: prod.dim_Dimname(Tensor self, Dimname dim, bool keepdim=False, *, ScalarType? dtype=None) -> Tensor	3437	- func: prod.dim_Dimname(Tensor self, Dimname dim, bool keepdim=False, *, ScalarType? dtype=None) -> Tensor
3428	variants: function, method	3438	variants: function, method
3429	supports_named_tensor: True	3439	supports_named_tensor: True
3430	npu_dispatch:	3440	npu_dispatch:
3431	NPU: prod_npu	3441	NPU: prod_npu
3432	#NPU: prod_npu_ext	3442	#NPU: prod_npu_ext
3433	-	3443	+
3434	- func: prod.Dimname_out(Tensor self, Dimname dim, bool keepdim=False, *, ScalarType? dtype=None, Tensor(a!) out) -> Tensor(a!)	3444	- func: prod.Dimname_out(Tensor self, Dimname dim, bool keepdim=False, *, ScalarType? dtype=None, Tensor(a!) out) -> Tensor(a!)
3435	supports_named_tensor: True	3445	supports_named_tensor: True
3436	npu_dispatch:	3446	npu_dispatch:
	@@ -3799,7 +3809,7 @@
3799		3809
3800	- func: norm_except_dim(Tensor v, int pow=2, int dim=0) -> Tensor	3810	- func: norm_except_dim(Tensor v, int pow=2, int dim=0) -> Tensor
3801	variants: function	3811	variants: function
3802	-	3812	+
3803	# VariableType::_weight_norm does not want to be given a gap in the autograd graph,	3813	# VariableType::_weight_norm does not want to be given a gap in the autograd graph,
3804	# so we don't define "dispatch" variants for it.	3814	# so we don't define "dispatch" variants for it.
3805	- func: _weight_norm(Tensor v, Tensor g, int dim=0) -> Tensor	3815	- func: _weight_norm(Tensor v, Tensor g, int dim=0) -> Tensor
	@@ -5005,12 +5015,12 @@
5005	variants: method, function	5015	variants: method, function
5006	npu_dispatch:	5016	npu_dispatch:
5007	NPU: bitwise_and_npu	5017	NPU: bitwise_and_npu
5008	-	5018	+
5009	- func: bitwise_and_.Scalar(Tensor(a!) self, Scalar other) -> Tensor(a!)	5019	- func: bitwise_and_.Scalar(Tensor(a!) self, Scalar other) -> Tensor(a!)
5010	variants: method	5020	variants: method
5011	npu_dispatch:	5021	npu_dispatch:
5012	NPU: bitwise_and_npu_	5022	NPU: bitwise_and_npu_
5013	-	5023	+
5014	- func: bitwise_and_.Tensor(Tensor(a!) self, Tensor other) -> Tensor(a!)	5024	- func: bitwise_and_.Tensor(Tensor(a!) self, Tensor other) -> Tensor(a!)
5015	variants: method	5025	variants: method
5016	npu_dispatch:	5026	npu_dispatch:
	@@ -5315,14 +5325,14 @@
5315	CPU: legacy::cpu::_th_addbmm_	5325	CPU: legacy::cpu::_th_addbmm_
5316	CUDA: legacy::cuda::_th_addbmm_	5326	CUDA: legacy::cuda::_th_addbmm_
5317	npu_dispatch:	5327	npu_dispatch:
5318	- NPU: addbmm_npu_	5328	+ NPU: addbmm_npu_
5319		5329
5320	- func: addbmm.out(Tensor self, Tensor batch1, Tensor batch2, *, Scalar beta=1, Scalar alpha=1, Tensor(a!) out) -> Tensor(a!)	5330	- func: addbmm.out(Tensor self, Tensor batch1, Tensor batch2, *, Scalar beta=1, Scalar alpha=1, Tensor(a!) out) -> Tensor(a!)
5321	dispatch:	5331	dispatch:
5322	CPU: legacy::cpu::_th_addbmm_out	5332	CPU: legacy::cpu::_th_addbmm_out
5323	CUDA: legacy::cuda::_th_addbmm_out	5333	CUDA: legacy::cuda::_th_addbmm_out
5324	npu_dispatch:	5334	npu_dispatch:
5325	- NPU: addbmm_out_npu	5335	+ NPU: addbmm_out_npu
5326		5336
5327	- func: addbmm(Tensor self, Tensor batch1, Tensor batch2, *, Scalar beta=1, Scalar alpha=1) -> Tensor	5337	- func: addbmm(Tensor self, Tensor batch1, Tensor batch2, *, Scalar beta=1, Scalar alpha=1) -> Tensor
5328	use_c10_dispatcher: full	5338	use_c10_dispatcher: full
	@@ -5331,7 +5341,7 @@
5331	CPU: legacy::cpu::_th_addbmm	5341	CPU: legacy::cpu::_th_addbmm
5332	CUDA: legacy::cuda::_th_addbmm	5342	CUDA: legacy::cuda::_th_addbmm
5333	npu_dispatch:	5343	npu_dispatch:
5334	- NPU: addbmm_npu	5344	+ NPU: addbmm_npu
5335		5345
5336	- func: addcdiv_(Tensor(a!) self, Tensor tensor1, Tensor tensor2, *, Scalar value=1) -> Tensor(a!)	5346	- func: addcdiv_(Tensor(a!) self, Tensor tensor1, Tensor tensor2, *, Scalar value=1) -> Tensor(a!)
5337	variants: method	5347	variants: method
	@@ -5794,7 +5804,7 @@
5794	- func: gather.dimname_out(Tensor self, Dimname dim, Tensor index, *, bool sparse_grad=False, Tensor(a!) out) -> Tensor(a!)	5804	- func: gather.dimname_out(Tensor self, Dimname dim, Tensor index, *, bool sparse_grad=False, Tensor(a!) out) -> Tensor(a!)
5795	npu_dispatch:	5805	npu_dispatch:
5796	NPU: gather_out_npu	5806	NPU: gather_out_npu
5797	-	5807	+
5798	- func: gather.dimname(Tensor self, Dimname dim, Tensor index, *, bool sparse_grad=False) -> Tensor	5808	- func: gather.dimname(Tensor self, Dimname dim, Tensor index, *, bool sparse_grad=False) -> Tensor
5799	variants: method, function	5809	variants: method, function
5800	npu_dispatch:	5810	npu_dispatch:
	@@ -6752,8 +6762,6 @@
6752	dispatch:	6762	dispatch:
6753	CPU: multilabel_margin_loss_backward_cpu_out	6763	CPU: multilabel_margin_loss_backward_cpu_out
6754	CUDA: legacy::cuda::_thnn_multilabel_margin_loss_backward_out	6764	CUDA: legacy::cuda::_thnn_multilabel_margin_loss_backward_out
6755	- npu_dispatch:
6756	- NPU: multilabel_margin_loss_backward_npu_out
6757		6765
6758	- func: multilabel_margin_loss_backward(Tensor grad_output, Tensor self, Tensor target, int reduction, Tensor is_target) -> Tensor	6766	- func: multilabel_margin_loss_backward(Tensor grad_output, Tensor self, Tensor target, int reduction, Tensor is_target) -> Tensor
6759	use_c10_dispatcher: full	6767	use_c10_dispatcher: full
	@@ -6761,8 +6769,6 @@
6761	dispatch:	6769	dispatch:
6762	CPU: multilabel_margin_loss_backward_cpu	6770	CPU: multilabel_margin_loss_backward_cpu
6763	CUDA: legacy::cuda::_thnn_multilabel_margin_loss_backward	6771	CUDA: legacy::cuda::_thnn_multilabel_margin_loss_backward
6764	- npu_dispatch:
6765	- NPU: multilabel_margin_loss_backward_npu
6766		6772
6767	- func: nll_loss.out(Tensor self, Tensor target, Tensor? weight=None, int reduction=Mean, int ignore_index=-100, *, Tensor(a!) out) -> Tensor(a!)	6773	- func: nll_loss.out(Tensor self, Tensor target, Tensor? weight=None, int reduction=Mean, int ignore_index=-100, *, Tensor(a!) out) -> Tensor(a!)
6768	python_module: nn	6774	python_module: nn
	@@ -7757,7 +7763,7 @@
7757	CUDA: upsample_bicubic2d_out_cuda	7763	CUDA: upsample_bicubic2d_out_cuda
7758	npu_dispatch:	7764	npu_dispatch:
7759	NPU: upsample_bicubic2d_out_npu	7765	NPU: upsample_bicubic2d_out_npu
7760	-	7766	+
7761	- func: upsample_bicubic2d(Tensor self, int[2] output_size, bool align_corners, float? scales_h=None, float? scales_w=None) -> Tensor	7767	- func: upsample_bicubic2d(Tensor self, int[2] output_size, bool align_corners, float? scales_h=None, float? scales_w=None) -> Tensor
7762	python_module: nn	7768	python_module: nn
7763	dispatch:	7769	dispatch:
	@@ -7765,7 +7771,7 @@
7765	CUDA: upsample_bicubic2d_cuda	7771	CUDA: upsample_bicubic2d_cuda
7766	npu_dispatch:	7772	npu_dispatch:
7767	NPU: upsample_bicubic2d_npu	7773	NPU: upsample_bicubic2d_npu
7768	-	7774	+
7769	- func: upsample_bicubic2d_backward.grad_input(Tensor grad_output, int[2] output_size, int[4] input_size, bool align_corners, float? scales_h=None, float? scales_w=None, *, Tensor(a!) grad_input) -> Tensor(a!)	7775	- func: upsample_bicubic2d_backward.grad_input(Tensor grad_output, int[2] output_size, int[4] input_size, bool align_corners, float? scales_h=None, float? scales_w=None, *, Tensor(a!) grad_input) -> Tensor(a!)
7770	python_module: nn	7776	python_module: nn
7771	dispatch:	7777	dispatch:
	@@ -7787,24 +7793,32 @@
7787	dispatch:	7793	dispatch:
7788	CPU: upsample_trilinear3d_out_cpu	7794	CPU: upsample_trilinear3d_out_cpu
7789	CUDA: upsample_trilinear3d_out_cuda	7795	CUDA: upsample_trilinear3d_out_cuda
		7796	+ npu_dispatch:
		7797	+ NPU: upsample_trilinear3d_out_npu
7790		7798
7791	- func: upsample_trilinear3d(Tensor self, int[3] output_size, bool align_corners, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor	7799	- func: upsample_trilinear3d(Tensor self, int[3] output_size, bool align_corners, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor
7792	python_module: nn	7800	python_module: nn
7793	dispatch:	7801	dispatch:
7794	CPU: upsample_trilinear3d_cpu	7802	CPU: upsample_trilinear3d_cpu
7795	CUDA: upsample_trilinear3d_cuda	7803	CUDA: upsample_trilinear3d_cuda
		7804	+ npu_dispatch:
		7805	+ NPU: upsample_trilinear3d_npu
7796		7806
7797	- func: upsample_trilinear3d_backward.grad_input(Tensor grad_output, int[3] output_size, int[5] input_size, bool align_corners, float? scales_d=None, float? scales_h=None, float? scales_w=None, *, Tensor(a!) grad_input) -> Tensor(a!)	7807	- func: upsample_trilinear3d_backward.grad_input(Tensor grad_output, int[3] output_size, int[5] input_size, bool align_corners, float? scales_d=None, float? scales_h=None, float? scales_w=None, *, Tensor(a!) grad_input) -> Tensor(a!)
7798	python_module: nn	7808	python_module: nn
7799	dispatch:	7809	dispatch:
7800	CPU: upsample_trilinear3d_backward_out_cpu	7810	CPU: upsample_trilinear3d_backward_out_cpu
7801	CUDA: upsample_trilinear3d_backward_out_cuda	7811	CUDA: upsample_trilinear3d_backward_out_cuda
		7812	+ npu_dispatch:
		7813	+ NPU: upsample_trilinear3d_backward_out_npu
7802		7814
7803	- func: upsample_trilinear3d_backward(Tensor grad_output, int[3] output_size, int[5] input_size, bool align_corners, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor	7815	- func: upsample_trilinear3d_backward(Tensor grad_output, int[3] output_size, int[5] input_size, bool align_corners, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor
7804	python_module: nn	7816	python_module: nn
7805	dispatch:	7817	dispatch:
7806	CPU: upsample_trilinear3d_backward_cpu	7818	CPU: upsample_trilinear3d_backward_cpu
7807	CUDA: upsample_trilinear3d_backward_cuda	7819	CUDA: upsample_trilinear3d_backward_cuda
		7820	+ npu_dispatch:
		7821	+ NPU: upsample_trilinear3d_backward_npu
7808		7822
7809	- func: upsample_nearest1d.out(Tensor self, int[1] output_size, float? scales=None, *, Tensor(a!) out) -> Tensor(a!)	7823	- func: upsample_nearest1d.out(Tensor self, int[1] output_size, float? scales=None, *, Tensor(a!) out) -> Tensor(a!)
7810	python_module: nn	7824	python_module: nn
	@@ -7876,6 +7890,8 @@
7876	dispatch:	7890	dispatch:
7877	CPU: upsample_nearest3d_out_cpu	7891	CPU: upsample_nearest3d_out_cpu
7878	CUDA: upsample_nearest3d_out_cuda	7892	CUDA: upsample_nearest3d_out_cuda
		7893	+ npu_dispatch:
		7894	+ NPU: upsample_nearest3d_out_npu
7879		7895
7880	- func: upsample_nearest3d(Tensor self, int[3] output_size, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor	7896	- func: upsample_nearest3d(Tensor self, int[3] output_size, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor
7881	python_module: nn	7897	python_module: nn
	@@ -7883,18 +7899,24 @@
7883	CPU: upsample_nearest3d_cpu	7899	CPU: upsample_nearest3d_cpu
7884	CUDA: upsample_nearest3d_cuda	7900	CUDA: upsample_nearest3d_cuda
7885	QuantizedCPU: quantized_upsample_nearest3d_cpu	7901	QuantizedCPU: quantized_upsample_nearest3d_cpu
		7902	+ npu_dispatch:
		7903	+ NPU: upsample_nearest3d_npu
7886		7904
7887	- func: upsample_nearest3d_backward.grad_input(Tensor grad_output, int[3] output_size, int[5] input_size, float? scales_d=None, float? scales_h=None, float? scales_w=None, *, Tensor(a!) grad_input) -> Tensor(a!)	7905	- func: upsample_nearest3d_backward.grad_input(Tensor grad_output, int[3] output_size, int[5] input_size, float? scales_d=None, float? scales_h=None, float? scales_w=None, *, Tensor(a!) grad_input) -> Tensor(a!)
7888	python_module: nn	7906	python_module: nn
7889	dispatch:	7907	dispatch:
7890	CPU: upsample_nearest3d_backward_out_cpu	7908	CPU: upsample_nearest3d_backward_out_cpu
7891	CUDA: upsample_nearest3d_backward_out_cuda	7909	CUDA: upsample_nearest3d_backward_out_cuda
		7910	+ npu_dispatch:
		7911	+ NPU: upsample_nearest3d_backward_out_npu
7892		7912
7893	- func: upsample_nearest3d_backward(Tensor grad_output, int[3] output_size, int[5] input_size, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor	7913	- func: upsample_nearest3d_backward(Tensor grad_output, int[3] output_size, int[5] input_size, float? scales_d=None, float? scales_h=None, float? scales_w=None) -> Tensor
7894	python_module: nn	7914	python_module: nn
7895	dispatch:	7915	dispatch:
7896	CPU: upsample_nearest3d_backward_cpu	7916	CPU: upsample_nearest3d_backward_cpu
7897	CUDA: upsample_nearest3d_backward_cuda	7917	CUDA: upsample_nearest3d_backward_cuda
		7918	+ npu_dispatch:
		7919	+ NPU: upsample_nearest3d_backward_npu
7898		7920
7899	- func: sigmoid_backward.grad_input(Tensor grad_output, Tensor output, *, Tensor(a!) grad_input) -> Tensor(a!)	7921	- func: sigmoid_backward.grad_input(Tensor grad_output, Tensor output, *, Tensor(a!) grad_input) -> Tensor(a!)
7900	python_module: nn	7922	python_module: nn
	@@ -8333,7 +8355,7 @@
8333	npu_dispatch_only:	8355	npu_dispatch_only:
8334	NPU: ptiou_npu	8356	NPU: ptiou_npu
8335		8357
8336	-- func: npu_nms_with_mask(Tensor input, Scalar iou_threshold) -> (Tensor, Tensor, Tensor)	8358	+- func: npu_nms_with_mask(Tensor input, Scalar iou_threshold) -> (Tensor, Tensor, Tensor)
8337	variants: function	8359	variants: function
8338	npu_dispatch_only:	8360	npu_dispatch_only:
8339	NPU: nms_with_mask_npu	8361	NPU: nms_with_mask_npu
	@@ -8406,7 +8428,7 @@
8406	variants: function, method	8428	variants: function, method
8407	npu_dispatch_only:	8429	npu_dispatch_only:
8408	NPU: indexing_npu	8430	NPU: indexing_npu
8409	-	8431	+
8410	- func: npu_indexing.out(Tensor self, int[] begin, int[] end, int[] strides, *, Tensor(a!) out) -> Tensor(a!)	8432	- func: npu_indexing.out(Tensor self, int[] begin, int[] end, int[] strides, *, Tensor(a!) out) -> Tensor(a!)
8411	npu_dispatch_only:	8433	npu_dispatch_only:
8412	NPU: indexing_out_npu	8434	NPU: indexing_out_npu
	@@ -8435,7 +8457,7 @@
8435	- func: npu_apply_adam(Tensor(a!) var, Tensor(b!) m, Tensor(c!) v, Scalar beta1_power, Scalar beta2_power, Scalar lr, Scalar beta1, Scalar beta2, Scalar epsilon, Tensor grad, bool? use_locking, bool? use_nesterov) -> (Tensor(a!), Tensor(b!), Tensor(c!))	8457	- func: npu_apply_adam(Tensor(a!) var, Tensor(b!) m, Tensor(c!) v, Scalar beta1_power, Scalar beta2_power, Scalar lr, Scalar beta1, Scalar beta2, Scalar epsilon, Tensor grad, bool? use_locking, bool? use_nesterov) -> (Tensor(a!), Tensor(b!), Tensor(c!))
8436	npu_dispatch_only:	8458	npu_dispatch_only:
8437	NPU: apply_adam_npu	8459	NPU: apply_adam_npu
8438	-	8460	+
8439	- func: npu_layer_norm_eval(Tensor input, int[] normalized_shape, Tensor? weight=None, Tensor? bias=None, float eps=1e-05) -> Tensor	8461	- func: npu_layer_norm_eval(Tensor input, int[] normalized_shape, Tensor? weight=None, Tensor? bias=None, float eps=1e-05) -> Tensor
8440	npu_dispatch_only:	8462	npu_dispatch_only:
8441	NPU: layer_norm_eval_npu	8463	NPU: layer_norm_eval_npu
	@@ -8464,7 +8486,7 @@
8464	npu_dispatch_only:	8486	npu_dispatch_only:
8465	NPU: confusion_transpose_backward_npu	8487	NPU: confusion_transpose_backward_npu
8466		8488
8467	-- func: npu_bmmV2(Tensor self, Tensor mat2) -> Tensor	8489	+- func: npu_bmmV2(Tensor self, Tensor mat2, int[] output_sizes) -> Tensor
8468	variants: function, method	8490	variants: function, method
8469	npu_dispatch_only:	8491	npu_dispatch_only:
8470	NPU: bmm_v2_npu	8492	NPU: bmm_v2_npu
	@@ -8512,14 +8534,6 @@
8512	npu_dispatch_only:	8534	npu_dispatch_only:
8513	NPU: grid_assign_positive_npu	8535	NPU: grid_assign_positive_npu
8514		8536
8515	-- func: global_step_inc() -> ()
8516	- variants: function
8517	- use_c10_dispatcher: full
8518	-
8519	-- func: set_start_fuzz_compile_step(int step) -> ()
8520	- variants: function
8521	- use_c10_dispatcher: full
8522	-
8523	- func: npu_mish_backward(Tensor grad, Tensor input) -> Tensor	8537	- func: npu_mish_backward(Tensor grad, Tensor input) -> Tensor
8524	npu_dispatch_only:	8538	npu_dispatch_only:
8525	NPU: mish_backward_npu	8539	NPU: mish_backward_npu
	@@ -8527,4 +8541,9 @@
8527	- func: npu_normalize_batch(Tensor self, Tensor seq_len, int normalize_type=0) -> Tensor	8541	- func: npu_normalize_batch(Tensor self, Tensor seq_len, int normalize_type=0) -> Tensor
8528	variants: function, method	8542	variants: function, method
8529	npu_dispatch_only:	8543	npu_dispatch_only:
8530	- NPU: normalize_batch_npu	8544	+ NPU: normalize_batch_npu
		8545	+
		8546	+- func: npu_masked_fill_range(Tensor self, Tensor start, Tensor end, Tensor value, int axis=-1) -> Tensor
		8547	+ variants: function, method
		8548	+ npu_dispatch_only:
		8549	+ NPU: masked_fill_range_npu

Msrc/aten/src/ATen/native/npu/AcosKernelNpu.cpp+2-10

	@@ -32,13 +32,7 @@ Tensor& acos_out_npu(Tensor& result, const Tensor& self) {
32	}	32	}
33		33
34	Tensor acos_npu(const Tensor& self) {	34	Tensor acos_npu(const Tensor& self) {
35	- // calculate the output size	35	+ Tensor result = OpPreparation::ApplyTensor(self);
36	- auto outputSize = input_same_output_size(self);
37	-
38	- // construct the output tensor of the NPU
39	- Tensor result = at::empty_with_format(
40	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
41	-
42	// calculate the output result of the NPU	36	// calculate the output result of the NPU
43	acos_out_npu(result, self);	37	acos_out_npu(result, self);
44		38
	@@ -46,9 +40,7 @@ Tensor acos_npu(const Tensor& self) {
46	}	40	}
47		41
48	Tensor& acos_npu_(Tensor& self) {	42	Tensor& acos_npu_(Tensor& self) {
49	- SmallVector<Tensor, N> inputs = {self};	43	+ OpPreparation::CheckMemory({self}, {self});
50	- SmallVector<Tensor, N> outputs = {self};
51	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
52		44
53	if (!NpuUtils::check_match(&self)) {	45	if (!NpuUtils::check_match(&self)) {
54	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	46	Tensor contiguousSelf = NpuUtils::format_contiguous(self);

Msrc/aten/src/ATen/native/npu/AddbmmKernelNpu.cpp+3-9

	@@ -12,10 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/CalcuOpUtil.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
17	-#include "ATen/native/npu/utils/NpuUtils.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		16
20	namespace at {	17	namespace at {
21	namespace native {	18	namespace native {
	@@ -45,8 +42,7 @@ Tensor addbmm_npu(
45	// calculate the output size	42	// calculate the output size
46	auto outputSize = addbmm_npu_output_size(self, batch1, batch2, beta, alpha);	43	auto outputSize = addbmm_npu_output_size(self, batch1, batch2, beta, alpha);
47	// construct the output tensor of the NPU	44	// construct the output tensor of the NPU
48	- Tensor result = at::empty_with_format(	45	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
49	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
50	// calculate the output result of the NPU	46	// calculate the output result of the NPU
51	addbmm_out_npu(result, self, batch1, batch2, beta, alpha);	47	addbmm_out_npu(result, self, batch1, batch2, beta, alpha);
52	return result;	48	return result;
	@@ -58,9 +54,7 @@ Tensor& addbmm_npu_(
58	const Tensor& batch2,	54	const Tensor& batch2,
59	Scalar beta,	55	Scalar beta,
60	Scalar alpha) {	56	Scalar alpha) {
61	- SmallVector<Tensor, N> inputs = {self, batch1, batch2};	57	+ OpPreparation::CheckMemory({self, batch1, batch2}, {self});
62	- SmallVector<Tensor, N> outputs = {self};
63	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
64	if (!NpuUtils::check_match(&self)) {	58	if (!NpuUtils::check_match(&self)) {
65	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	59	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
66	Tensor result = addbmm_out_npu(contiguousSelf, contiguousSelf, batch1, batch2, beta, alpha);	60	Tensor result = addbmm_out_npu(contiguousSelf, contiguousSelf, batch1, batch2, beta, alpha);

Msrc/aten/src/ATen/native/npu/AddcdivKernelNpu.cpp+3-8

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -46,9 +45,7 @@ Tensor addcdiv_npu(
46	Scalar value) {	45	Scalar value) {
47	auto divOutputSize = broadcast_ops_npu_output_size(tensor1, tensor2);	46	auto divOutputSize = broadcast_ops_npu_output_size(tensor1, tensor2);
48	auto outputSize = broadcast_ops_npu_output_size(self.sizes(), divOutputSize);	47	auto outputSize = broadcast_ops_npu_output_size(self.sizes(), divOutputSize);
49	- Tensor result = at::empty_with_format(	48	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
50	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
51	-
52	addcdiv_out_npu(result, self, tensor1, tensor2, value);	49	addcdiv_out_npu(result, self, tensor1, tensor2, value);
53		50
54	return result;	51	return result;
	@@ -59,9 +56,7 @@ Tensor& addcdiv_npu_(
59	const Tensor& tensor1,	56	const Tensor& tensor1,
60	const Tensor& tensor2,	57	const Tensor& tensor2,
61	Scalar value) {	58	Scalar value) {
62	- SmallVector<Tensor, N> inputs = {self, tensor1, tensor2};	59	+ OpPreparation::CheckMemory({self, tensor1, tensor2}, {self});
63	- SmallVector<Tensor, N> outputs = {self};
64	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
65	if (!NpuUtils::check_match(&self)) {	60	if (!NpuUtils::check_match(&self)) {
66	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	61	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
67	Tensor result = addcdiv_out_npu(contiguousSelf, contiguousSelf, tensor1, tensor2, value);	62	Tensor result = addcdiv_out_npu(contiguousSelf, contiguousSelf, tensor1, tensor2, value);

Msrc/aten/src/ATen/native/npu/AddcmulKernelNpu.cpp+2-6

	@@ -50,8 +50,7 @@ Tensor& addcmul_out_npu(
50	OpPreparation::CheckOut(	50	OpPreparation::CheckOut(
51	{self},	51	{self},
52	result,	52	result,
53	- CalcuOpUtil::get_tensor_npu_format(self),	53	+ self,
54	- self.scalar_type(),
55	outputSize);	54	outputSize);
56		55
57	OpPipeWithDefinedOut pipe;	56	OpPipeWithDefinedOut pipe;
	@@ -81,10 +80,7 @@ Tensor& addcmul_npu_(
81	const Tensor& tensor1,	80	const Tensor& tensor1,
82	const Tensor& tensor2,	81	const Tensor& tensor2,
83	Scalar value) {	82	Scalar value) {
84	- SmallVector<Tensor, N> inputs = {self};	83	+ OpPreparation::CheckMemory({self}, {self});
85	- SmallVector<Tensor, N> outputs = {self};
86	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
87	-
88	if (!NpuUtils::check_match(&self)) {	84	if (!NpuUtils::check_match(&self)) {
89	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	85	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
90	Tensor result = addcmul_out_npu_nocheck(	86	Tensor result = addcmul_out_npu_nocheck(

Msrc/aten/src/ATen/native/npu/AddmvKernelNpu.cpp+3-12

	@@ -12,8 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/CalcuOpUtil.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
17	#include "ATen/native/npu/utils/NpuUtils.h"	16	#include "ATen/native/npu/utils/NpuUtils.h"
18		17
19	namespace at {	18	namespace at {
	@@ -64,14 +63,8 @@ Tensor addmv_npu(
64	Scalar alpha) {	63	Scalar alpha) {
65		64
66	check_1d(vec, "vec", "addmv");	65	check_1d(vec, "vec", "addmv");
67	- // calculate the output size
68	auto outputSize = addmv_npu_output_size(self, mat, vec, beta, alpha);	66	auto outputSize = addmv_npu_output_size(self, mat, vec, beta, alpha);
69	-	67	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
70	- // construct the output tensor of the NPU
71	- Tensor result = at::empty_with_format(
72	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
73	-
74	- // calculate the output result of the NPU
75	addmv_out_npu(result, self, mat, vec, beta, alpha);	68	addmv_out_npu(result, self, mat, vec, beta, alpha);
76		69
77	return result;	70	return result;
	@@ -85,9 +78,7 @@ Tensor& addmv_npu_(
85	Scalar alpha) {	78	Scalar alpha) {
86		79
87	check_1d(vec, "vec", "addmv");	80	check_1d(vec, "vec", "addmv");
88	- SmallVector<Tensor, N> inputs = {self, mat, vec};	81	+ OpPreparation::CheckMemory({self, mat, vec}, {self});
89	- SmallVector<Tensor, N> outputs = {self};
90	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
91	if (!NpuUtils::check_match(&self)) {	82	if (!NpuUtils::check_match(&self)) {
92	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	83	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
93	Tensor result =	84	Tensor result =

Msrc/aten/src/ATen/native/npu/AddrKernelNpu.cpp+3-13

	@@ -12,8 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/CalcuOpUtil.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
17	#include "ATen/native/npu/utils/NpuUtils.h"	16	#include "ATen/native/npu/utils/NpuUtils.h"
18		17
19	namespace at {	18	namespace at {
	@@ -64,15 +63,8 @@ Tensor _addr_npu(
64	const Tensor& vec2,	63	const Tensor& vec2,
65	Scalar beta,	64	Scalar beta,
66	Scalar alpha) {	65	Scalar alpha) {
67	-
68	- // calculate the output size
69	auto outputSize = addr_npu_output_size(self, vec1, vec2, beta, alpha);	66	auto outputSize = addr_npu_output_size(self, vec1, vec2, beta, alpha);
70	-	67	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
71	- // construct the output tensor of the NPU
72	- Tensor result = at::empty_with_format(
73	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
74	-
75	- // calculate the output result of the NPU
76	_addr_out_npu(result, self, vec1, vec2, beta, alpha);	68	_addr_out_npu(result, self, vec1, vec2, beta, alpha);
77		69
78	return result;	70	return result;
	@@ -95,9 +87,7 @@ Tensor& _addr_npu_(
95	const Tensor& vec2,	87	const Tensor& vec2,
96	Scalar beta,	88	Scalar beta,
97	Scalar alpha) {	89	Scalar alpha) {
98	- SmallVector<Tensor, N> inputs = {self, vec1, vec2};	90	+ OpPreparation::CheckMemory({self, vec1, vec2}, {self});
99	- SmallVector<Tensor, N> outputs = {self};
100	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
101	if (!NpuUtils::check_match(&self)) {	91	if (!NpuUtils::check_match(&self)) {
102	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	92	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
103	Tensor result =	93	Tensor result =

Msrc/aten/src/ATen/native/npu/AffineGridGeneratorBackwardKernelNpu.cpp+1-2

	@@ -35,8 +35,7 @@ Tensor& affine_grid_generator_backward_nocheck(
35	const Tensor& grad,	35	const Tensor& grad,
36	IntArrayRef size,	36	IntArrayRef size,
37	bool align_corners) {	37	bool align_corners) {
38	- Tensor assist = at::empty_with_format(	38	+ Tensor assist = OpPreparation::ApplyTensor(grad, {size[0], size[2], size[3], 3});
39	- {size[0], size[2], size[3], 3}, grad.options(), CalcuOpUtil::get_tensor_npu_format(grad));
40	assist.select(-1, 0).copy_(_linspace_from_neg_one(grad, size[3], align_corners));	39	assist.select(-1, 0).copy_(_linspace_from_neg_one(grad, size[3], align_corners));
41	assist.select(-1, 1).copy_(_linspace_from_neg_one(grad, size[2], align_corners).unsqueeze_(-1));	40	assist.select(-1, 1).copy_(_linspace_from_neg_one(grad, size[2], align_corners).unsqueeze_(-1));
42	assist.select(-1, 2).fill_(1);	41	assist.select(-1, 2).fill_(1);

Msrc/aten/src/ATen/native/npu/AffineGridGeneratorKernelNpu.cpp+2-4

	@@ -17,7 +17,7 @@
17	#include "ATen/native/npu/utils/CalcuOpUtil.h"	17	#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	18	#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
19	#include "ATen/native/npu/utils/NpuUtils.h"	19	#include "ATen/native/npu/utils/NpuUtils.h"
20	-#include "ATen/native/npu/utils/OpTemplate.h"	20	+#include "ATen/native/npu/utils/OpAdapter.h"
21		21
22	namespace at {	22	namespace at {
23	namespace native {	23	namespace native {
	@@ -62,9 +62,7 @@ Tensor affine_grid_generator_npu(
62	}	62	}
63		63
64	// construct the output tensor of the NPU	64	// construct the output tensor of the NPU
65	- Tensor result = at::empty_with_format(	65	+ Tensor result = OpPreparation::ApplyTensor(theta, outputSize);
66	- outputSize, theta.options(), CalcuOpUtil::get_tensor_npu_format(theta));
67	-
68	// calculate the output result of the NPU	66	// calculate the output result of the NPU
69	affine_grid_generator_npu_nocheck(	67	affine_grid_generator_npu_nocheck(
70	result,	68	result,

Msrc/aten/src/ATen/native/npu/AllKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	#include "c10/npu/OptionsManager.h"	19	#include "c10/npu/OptionsManager.h"
19		20
20	namespace at {	21	namespace at {

Msrc/aten/src/ATen/native/npu/AnyKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/ArangeKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/ArgminKernelNpu.cpp+2-2

	@@ -32,10 +32,10 @@ Tensor argmin_npu(const Tensor& self, optional<int64_t> dim, bool keepdim) {
32	auto outputSize = reduce_ops_npu_output_size(input, realDim, realKeepDim);	32	auto outputSize = reduce_ops_npu_output_size(input, realDim, realKeepDim);
33		33
34	// construct the output tensor of the NPU	34	// construct the output tensor of the NPU
35	- Tensor result = at::empty_with_format(	35	+ Tensor result = OpPreparation::ApplyTensor(
36	outputSize,	36	outputSize,
37	self.options().dtype(at::kInt),	37	self.options().dtype(at::kInt),
38	- CalcuOpUtil::get_tensor_npu_format(self));	38	+ self);
39	SmallVector<int64_t, N> DimVec = {realDim};	39	SmallVector<int64_t, N> DimVec = {realDim};
40	// calculate the output result of the NPU	40	// calculate the output result of the NPU
41	OpCommand cmd;	41	OpCommand cmd;

Msrc/aten/src/ATen/native/npu/ArgsortKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include<ATen/NamedTensorUtils.h>
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/AsinKernelNpu.cpp+3-17

	@@ -12,8 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"
17		16
18	namespace at {	17	namespace at {
19	namespace native {	18	namespace native {
	@@ -32,26 +31,13 @@ Tensor& asin_out_npu(
32	}	31	}
33		32
34	Tensor asin_npu(const Tensor& self) {	33	Tensor asin_npu(const Tensor& self) {
35	- // calculate the output size	34	+ Tensor result = OpPreparation::ApplyTensor(self);
36	- auto outputSize = input_same_output_size(self);
37	-
38	- // construct the output tensor of the NPU
39	- Tensor result = at::empty_with_format(
40	- outputSize,
41	- self.options(),
42	- CalcuOpUtil::get_tensor_npu_format(self));
43	-
44	- // calculate the output result of the NPU
45	asin_out_npu(result, self);	35	asin_out_npu(result, self);
46	-
47	return result;	36	return result;
48	}	37	}
49		38
50	Tensor& asin_npu_(Tensor& self) {	39	Tensor& asin_npu_(Tensor& self) {
51	- SmallVector<Tensor, N> inputs = {self};	40	+ OpPreparation::CheckMemory({self}, {self});
52	- SmallVector<Tensor, N> outputs = {self};
53	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
54	-
55	if (!NpuUtils::check_match(&self)) {	41	if (!NpuUtils::check_match(&self)) {
56	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	42	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
57	Tensor result = asin_out_npu(contiguousSelf, contiguousSelf);	43	Tensor result = asin_out_npu(contiguousSelf, contiguousSelf);

Msrc/aten/src/ATen/native/npu/Atan2KernelNpu.cpp+4-17

	@@ -12,8 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"
17		16
18	namespace at {	17	namespace at {
19	namespace native {	18	namespace native {
	@@ -44,8 +43,7 @@ Tensor& atan2_out_npu(
44	OpPreparation::CheckOut(	43	OpPreparation::CheckOut(
45	{self},	44	{self},
46	result,	45	result,
47	- CalcuOpUtil::get_tensor_npu_format(self),	46	+ self,
48	- self.scalar_type(),
49	outputSize);	47	outputSize);
50		48
51	atan2_out_npu_nocheck(result, self, other);	49	atan2_out_npu_nocheck(result, self, other);
	@@ -54,25 +52,14 @@ Tensor& atan2_out_npu(
54	}	52	}
55		53
56	Tensor atan2_npu(const Tensor& self, const Tensor& other) {	54	Tensor atan2_npu(const Tensor& self, const Tensor& other) {
57	- // calculate the output size
58	auto outputSize = broadcast_ops_npu_output_size(self, other);	55	auto outputSize = broadcast_ops_npu_output_size(self, other);
59	-	56	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
60	- // construct the output tensor of the NPU
61	- Tensor result = at::empty_with_format(
62	- outputSize,
63	- self.options(),
64	- CalcuOpUtil::get_tensor_npu_format(self));
65	-
66	- // calculate the output result of the NPU
67	atan2_out_npu_nocheck(result, self, other);	57	atan2_out_npu_nocheck(result, self, other);
68	-
69	return result;	58	return result;
70	}	59	}
71		60
72	Tensor& atan2_npu_(Tensor& self, const Tensor& other) {	61	Tensor& atan2_npu_(Tensor& self, const Tensor& other) {
73	- SmallVector<Tensor, N> inputs = {self, other};	62	+ OpPreparation::CheckMemory({self, other}, {self});
74	- SmallVector<Tensor, N> outputs = {self};
75	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
76		63
77	if (!NpuUtils::check_match(&self)) {	64	if (!NpuUtils::check_match(&self)) {
78	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	65	Tensor contiguousSelf = NpuUtils::format_contiguous(self);

Msrc/aten/src/ATen/native/npu/AtanKernelNpu.cpp+2-11

	@@ -31,23 +31,14 @@ Tensor& atan_out_npu(Tensor& result, const Tensor& self) {
31	}	31	}
32		32
33	Tensor atan_npu(const Tensor& self) {	33	Tensor atan_npu(const Tensor& self) {
34	- //calculate the output size	34	+ Tensor result = OpPreparation::ApplyTensor(self);
35	- auto outputSize = input_same_output_size(self);
36	-
37	- //construct the output tensor of the NPU
38	- Tensor result = at::empty_with_format(
39	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
40	-
41	//calculate the output result of the NPU	35	//calculate the output result of the NPU
42	atan_out_npu(result, self);	36	atan_out_npu(result, self);
43	return result;	37	return result;
44	}	38	}
45		39
46	Tensor& atan_npu_(Tensor& self) {	40	Tensor& atan_npu_(Tensor& self) {
47	- SmallVector<Tensor, N> inputs = {self};	41	+ OpPreparation::CheckMemory({self}, {self});
48	- SmallVector<Tensor, N> outputs = {self};
49	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
50	-
51	if (!NpuUtils::check_match(&self)) {	42	if (!NpuUtils::check_match(&self)) {
52	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	43	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
53	Tensor result = atan_out_npu(contiguousSelf, contiguousSelf);	44	Tensor result = atan_out_npu(contiguousSelf, contiguousSelf);

Msrc/aten/src/ATen/native/npu/BatchNMSKernelNpu.cpp+8-9

	@@ -28,25 +28,24 @@ std::tuple<Tensor, Tensor, Tensor, Tensor> batch_nms_npu(
28	bool change_coordinate_frame,	28	bool change_coordinate_frame,
29	bool transpose_box) {	29	bool transpose_box) {
30	// construct the output tensor of the NPU	30	// construct the output tensor of the NPU
31	- Tensor nmsed_boxes = at::empty_with_format(	31	+ Tensor nmsed_boxes = OpPreparation::ApplyTensor(
32	{self.size(0), max_total_size, 4},	32	{self.size(0), max_total_size, 4},
33	self.options().dtype(at::kHalf),	33	self.options().dtype(at::kHalf),
34	- CalcuOpUtil::get_tensor_npu_format(self));	34	+ self);
35	-	35	+ Tensor nmsed_scores = OpPreparation::ApplyTensor(
36	- Tensor nmsed_scores = at::empty_with_format(
37	{self.size(0), max_total_size},	36	{self.size(0), max_total_size},
38	self.options().dtype(at::kHalf),	37	self.options().dtype(at::kHalf),
39	- CalcuOpUtil::get_tensor_npu_format(self));	38	+ self);
40		39
41	- Tensor nmsed_classes = at::empty_with_format(	40	+ Tensor nmsed_classes = OpPreparation::ApplyTensor(
42	{self.size(0), max_total_size},	41	{self.size(0), max_total_size},
43	self.options().dtype(at::kHalf),	42	self.options().dtype(at::kHalf),
44	- CalcuOpUtil::get_tensor_npu_format(self));	43	+ self);
45		44
46	- Tensor nmsed_num = at::empty_with_format(	45	+ Tensor nmsed_num = OpPreparation::ApplyTensor(
47	{self.size(0)},	46	{self.size(0)},
48	self.options().dtype(at::kInt),	47	self.options().dtype(at::kInt),
49	- CalcuOpUtil::get_tensor_npu_format(self));	48	+ self);
50		49
51	OpCommand cmd;	50	OpCommand cmd;
52	cmd.Name("BatchMultiClassNonMaxSuppression")	51	cmd.Name("BatchMultiClassNonMaxSuppression")

Msrc/aten/src/ATen/native/npu/BernoulliKernelNpu.cpp+3-10

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -44,10 +43,7 @@ Tensor& bernoulli_out_npu(Tensor& result, const Tensor& self, const Tensor& p) {
44	}	43	}
45		44
46	Tensor& bernoulli_npu_(Tensor& self, double p, Generator* gen) {	45	Tensor& bernoulli_npu_(Tensor& self, double p, Generator* gen) {
47	- SmallVector<Tensor, N> inputs = {self};	46	+ OpPreparation::CheckMemory({self}, {self});
48	- SmallVector<Tensor, N> outputs = {self};
49	-
50	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
51	ScalarType selfType = self.scalar_type();	47	ScalarType selfType = self.scalar_type();
52	Tensor selfFp32 = self;	48	Tensor selfFp32 = self;
53	if (self.scalar_type() == ScalarType::Half) {	49	if (self.scalar_type() == ScalarType::Half) {
	@@ -70,10 +66,7 @@ Tensor& bernoulli_npu_(Tensor& self, double p, Generator* gen) {
70	}	66	}
71		67
72	Tensor& bernoulli_npu_(Tensor& self, const Tensor& p, Generator* gen) {	68	Tensor& bernoulli_npu_(Tensor& self, const Tensor& p, Generator* gen) {
73	- SmallVector<Tensor, N> inputs = {self};	69	+ OpPreparation::CheckMemory({self}, {self});
74	- SmallVector<Tensor, N> outputs = {self};
75	-
76	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
77	ScalarType selfType = self.scalar_type();	70	ScalarType selfType = self.scalar_type();
78	Tensor selfFp32 = self;	71	Tensor selfFp32 = self;
79	Tensor pFp32 = OpPreparation::CastBackToOriFormat(p);;	72	Tensor pFp32 = OpPreparation::CastBackToOriFormat(p);;

Msrc/aten/src/ATen/native/npu/BinaryCrossEntropyBackwardKernelNpu.cpp+2-9

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -62,13 +61,7 @@ Tensor binary_cross_entropy_backward_npu(
62	const Tensor& target,	61	const Tensor& target,
63	const Tensor& weight,	62	const Tensor& weight,
64	int64_t reduction) {	63	int64_t reduction) {
65	- // calculate the output size	64	+ Tensor gradInput = OpPreparation::ApplyTensor(self);
66	- auto outputSize = input_same_output_size(self);
67	-
68	- // construct the output tensor of the NPU
69	- Tensor gradInput = at::empty_with_format(
70	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
71	-
72	// calculate the output result of the NPU	65	// calculate the output result of the NPU
73	binary_cross_entropy_backward_out_npu(	66	binary_cross_entropy_backward_out_npu(
74	gradInput, grad_output, self, target, weight, reduction);	67	gradInput, grad_output, self, target, weight, reduction);

Msrc/aten/src/ATen/native/npu/BinaryCrossEntropyKernelNpu.cpp+2-6

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -68,10 +67,7 @@ Tensor binary_cross_entropy_npu(
68	}	67	}
69		68
70	// construct the output tensor of the NPU	69	// construct the output tensor of the NPU
71	- Tensor result = at::empty_with_format(	70	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
72	- outputSize,
73	- self.options(),
74	- CalcuOpUtil::get_tensor_npu_format(self));
75		71
76	// calculate the output result of the NPU	72	// calculate the output result of the NPU
77	binary_cross_entropy_out_npu(result, self, target, weight, reduction);	73	binary_cross_entropy_out_npu(result, self, target, weight, reduction);

Msrc/aten/src/ATen/native/npu/BinaryCrossEntropyWithLogitsBackwardKernelNpu.cpp+2-8

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -28,12 +27,7 @@ Tensor binary_cross_entropy_with_logits_backward_npu(
28	const Tensor& weight,	27	const Tensor& weight,
29	const Tensor& pos_weight,	28	const Tensor& pos_weight,
30	int64_t reduction) {	29	int64_t reduction) {
31	- // calculate the output size	30	+ Tensor gradInput = OpPreparation::ApplyTensor(self);
32	- auto outputSize = input_same_output_size(self);
33	-
34	- // construct the output tensor of the NPU
35	- Tensor gradInput = at::empty_with_format(
36	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
37		31
38	// calculate the output result of the NPU	32	// calculate the output result of the NPU
39	Tensor weightTensor;	33	Tensor weightTensor;

Msrc/aten/src/ATen/native/npu/BinaryCrossEntropyWithLogitsKernelNpu.cpp+2-2

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19		19
20	namespace at {	20	namespace at {
21	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/BitwiseAndKernelNpu.cpp+2-2

	@@ -12,8 +12,8 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"	16	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
17		17
18	namespace at {	18	namespace at {
19	namespace native {	19	namespace native {

Msrc/aten/src/ATen/native/npu/BitwiseNotKernelNpu.cpp+1-3

	@@ -38,9 +38,7 @@ Tensor& bitwise_not_out_npu(Tensor& result, const Tensor& self) {
38	OpPreparation::CheckOut(	38	OpPreparation::CheckOut(
39	{self},	39	{self},
40	result,	40	result,
41	- CalcuOpUtil::get_tensor_npu_format(self),	41	+ self);
42	- self.scalar_type(),
43	- self.sizes());
44		42
45	OpPipeWithDefinedOut pipe;	43	OpPipeWithDefinedOut pipe;
46	return pipe.CheckMemory({self}, {result})	44	return pipe.CheckMemory({self}, {result})

Msrc/aten/src/ATen/native/npu/BitwiseOrKernelNpu.cpp+6-19

	@@ -12,8 +12,8 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"	16	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
17		17
18	namespace at {	18	namespace at {
19	namespace native {	19	namespace native {
	@@ -44,9 +44,7 @@ Tensor& bitwise_or_out_npu(
44	OpPreparation::CheckOut(	44	OpPreparation::CheckOut(
45	{self},	45	{self},
46	result,	46	result,
47	- CalcuOpUtil::get_tensor_npu_format(self),	47	+ self);
48	- self.scalar_type(),
49	- self.sizes());
50		48
51	bitwise_or_out_npu_nocheck(result, self, other);	49	bitwise_or_out_npu_nocheck(result, self, other);
52		50
	@@ -120,11 +118,7 @@ Tensor bitwise_or_npu(const Tensor& self, const Tensor& other) {
120	auto outputSize = broadcast_ops_npu_output_size(self, other);	118	auto outputSize = broadcast_ops_npu_output_size(self, other);
121		119
122	// construct the output tensor of the NPU	120	// construct the output tensor of the NPU
123	- Tensor result = at::empty_with_format(	121	+ Tensor result = OpPreparation::ApplyTensor(outputTensor, outputSize);
124	- outputSize,
125	- outputTensor.options(),
126	- CalcuOpUtil::get_tensor_npu_format(outputTensor));
127	-
128	// calculate the output result of the NPU	122	// calculate the output result of the NPU
129	bitwise_or_out_npu_nocheck(result, self, other);	123	bitwise_or_out_npu_nocheck(result, self, other);
130		124
	@@ -132,12 +126,7 @@ Tensor bitwise_or_npu(const Tensor& self, const Tensor& other) {
132	}	126	}
133		127
134	Tensor bitwise_or_npu(const Tensor& self, Scalar other) {	128	Tensor bitwise_or_npu(const Tensor& self, Scalar other) {
135	- // calculate the output size	129	+ Tensor result = OpPreparation::ApplyTensor(self);
136	- auto outputSize = input_same_output_size(self);
137	-
138	- // construct the output tensor of the NPU
139	- Tensor result = at::empty_with_format(
140	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
141		130
142	// calculate the output result of the NPU	131	// calculate the output result of the NPU
143	bitwise_or_out_npu_nocheck(result, self, other);	132	bitwise_or_out_npu_nocheck(result, self, other);
	@@ -146,9 +135,7 @@ Tensor bitwise_or_npu(const Tensor& self, Scalar other) {
146	}	135	}
147		136
148	Tensor& bitwise_or_npu_(Tensor& self, const Tensor& other) {	137	Tensor& bitwise_or_npu_(Tensor& self, const Tensor& other) {
149	- SmallVector<Tensor, N> inputs = {self, other};	138	+ OpPreparation::CheckMemory({self, other}, {self});
150	- SmallVector<Tensor, N> outputs = {self};
151	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
152		139
153	if (!NpuUtils::check_match(&self)) {	140	if (!NpuUtils::check_match(&self)) {
154	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	141	Tensor contiguousSelf = NpuUtils::format_contiguous(self);

Msrc/aten/src/ATen/native/npu/BitwiseXorKernelNpu.cpp+6-23

	@@ -12,8 +12,8 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"	16	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
17		17
18	namespace at {	18	namespace at {
19	namespace native {	19	namespace native {
	@@ -44,9 +44,7 @@ Tensor& bitwise_xor_out_npu(
44	OpPreparation::CheckOut(	44	OpPreparation::CheckOut(
45	{self},	45	{self},
46	result,	46	result,
47	- CalcuOpUtil::get_tensor_npu_format(self),	47	+ self);
48	- self.scalar_type(),
49	- self.sizes());
50		48
51	bitwise_xor_out_npu_nocheck(result, self, other);	49	bitwise_xor_out_npu_nocheck(result, self, other);
52		50
	@@ -122,11 +120,7 @@ Tensor bitwise_xor_npu(const Tensor& self, const Tensor& other) {
122	auto outputSize = broadcast_ops_npu_output_size(self, other);	120	auto outputSize = broadcast_ops_npu_output_size(self, other);
123		121
124	// construct the output tensor of the NPU	122	// construct the output tensor of the NPU
125	- Tensor result = at::empty_with_format(	123	+ Tensor result = OpPreparation::ApplyTensor(outputTensor, outputSize);
126	- outputSize,
127	- outputTensor.options(),
128	- CalcuOpUtil::get_tensor_npu_format(outputTensor));
129	-
130	// calculate the output result of the NPU	124	// calculate the output result of the NPU
131	bitwise_xor_out_npu_nocheck(result, self, other);	125	bitwise_xor_out_npu_nocheck(result, self, other);
132		126
	@@ -134,15 +128,7 @@ Tensor bitwise_xor_npu(const Tensor& self, const Tensor& other) {
134	}	128	}
135		129
136	Tensor bitwise_xor_npu(const Tensor& self, Scalar other) {	130	Tensor bitwise_xor_npu(const Tensor& self, Scalar other) {
137	- // calculate the output size	131	+ Tensor result = OpPreparation::ApplyTensor(self);
138	- auto outputSize = input_same_output_size(self);
139	-
140	- // construct the output tensor of the NPU
141	- Tensor result = at::empty_with_format(
142	- outputSize,
143	- self.options(),
144	- CalcuOpUtil::get_tensor_npu_format(self));
145	-
146	// calculate the output result of the NPU	132	// calculate the output result of the NPU
147	bitwise_xor_out_npu_nocheck(result, self, other);	133	bitwise_xor_out_npu_nocheck(result, self, other);
148		134
	@@ -150,10 +136,7 @@ Tensor bitwise_xor_npu(const Tensor& self, Scalar other) {
150	}	136	}
151		137
152	Tensor& bitwise_xor_npu_(Tensor& self, const Tensor& other) {	138	Tensor& bitwise_xor_npu_(Tensor& self, const Tensor& other) {
153	- SmallVector<Tensor, N> inputs = {self, other};	139	+ OpPreparation::CheckMemory({self, other}, {self});
154	- SmallVector<Tensor, N> outputs = {self};
155	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
156	-
157	if (!NpuUtils::check_match(&self)) {	140	if (!NpuUtils::check_match(&self)) {
158	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	141	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
159	Tensor result = bitwise_xor_out_npu_nocheck(contiguousSelf, contiguousSelf, other);	142	Tensor result = bitwise_xor_out_npu_nocheck(contiguousSelf, contiguousSelf, other);

Msrc/aten/src/ATen/native/npu/BmmKernelNpu.cpp+2-2

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19	#include "c10/npu/OptionsManager.h"	19	#include "c10/npu/OptionsManager.h"
20		20
21	namespace at {	21	namespace at {

Msrc/aten/src/ATen/native/npu/BmmV2KernelNpu.cpp+244-38

 // See the License for the specific language governing permissions and
 // limitations under the License.
++#include "ATen/native/npu/utils/OpAdapter.h"
++#include "ATen/native/npu/utils/CalcuOpUtil.h"
 namespace at {
 namespace native {
 using namespace at::native::npu;
++bool is_transpose_last_two_dims_v2(const Tensor& tensor) {
++  if (tensor.dim() < 2) {
++    return false;
++  }
++  auto storage_size = tensor.storage().get_npu_desc().storage_sizes_;
++  int64_t numel = at::prod_intlist(storage_size);
++
++  int64_t dim1 = tensor.dim() - 1;
++  int64_t dim2 = tensor.dim() - 2;
++
++  auto tensor_desc = tensor.storage().get_npu_desc();
++  if (tensor_desc.base_sizes_.size() == tensor.dim() &&
++      tensor.stride(dim2) == 1 && tensor.stride(dim1) == tensor.size(dim2) &&
++      tensor.size(dim1) == tensor_desc.base_sizes_[dim2] &&
++      tensor.size(dim2) == tensor_desc.base_sizes_[dim1] &&
++      tensor.storage().size() == numel) {
++    return true;
   } else {
++    return false;
+-    IntArrayRef batch_tensor1(mat1.sizes().data(), std::max<int64_t>(dim_tensor1 - 2, 0));
+-    SmallVector<int64_t, SIZE> output_size = array_to_small_vector(batch_tensor1);
+-    output_size.emplace_back(mat1.size(-2));
+-    output_size.emplace_back(mat2.size(-1));
+-    return output_size;
   }
 }
++SmallVector<int64_t, SIZE> bmm_v2_output_size(const Tensor& mat1, const Tensor& mat2) {
++  auto dim_tensor1 = mat1.dim();
++  auto dim_tensor2 = mat2.dim();
++  int64_t m = dim_tensor1 == 1 ? 1 : mat1.size(-2);
++  int64_t n = dim_tensor2 == 1 ? 1 : mat2.size(-1);
+-	Tensor result;
++  auto batch_a = array_to_small_vector(IntArrayRef(mat1.sizes().data(), std::max<int64_t>(dim_tensor1 - 2, 0)));
++  auto batch_b = array_to_small_vector(IntArrayRef(mat2.sizes().data(), std::max<int64_t>(dim_tensor2 - 2, 0)));
++
++  batch_a.insert(batch_a.begin(), std::max<int64_t>(batch_a.size(), batch_b.size()) - batch_a.size(), 1);
++  batch_b.insert(batch_b.begin(), std::max<int64_t>(batch_a.size(), batch_b.size()) - batch_b.size(), 1);
++
++  SmallVector<int64_t, SIZE> output_size;
++  for (size_t i = 0; i < batch_a.size(); ++i) {
++    if (batch_a[i] == 1) {
++      output_size.emplace_back(batch_b[i]);
++    } else if (batch_b[i] == 1) {
++      output_size.emplace_back(batch_a[i]);
++    } else if (batch_a[i] != batch_b[i]) {
++      AT_ERROR("mat1 and mat2 cannot broadcast, but they are mat1 ",
++          mat1.sizes().data(), " mat2 ", mat2.sizes().data());
++    } else {
++      output_size.emplace_back(batch_a[i]);
++    }
++  }
++  output_size.emplace_back(m);
++  output_size.emplace_back(n);
++
++  return output_size;
++}
++
++Tensor pure_bmm_v2_npu(const Tensor& self, const Tensor& mat2, const SmallVector<int64_t, SIZE>& output_size) {
++  auto tensor1 = self.dim() == 1 ? self.view({1, self.size(0)}) : self;
++  auto tensor2 = mat2.dim() == 1 ? mat2.view({mat2.size(0), 1}) : mat2;
++
++  Tensor result;
++
++  if ((tensor1.scalar_type() == ScalarType::Half)) {
++    result = at::empty_with_format(output_size, tensor1.options(), ACL_FORMAT_FRACTAL_NZ);
   } else {
++    result = at::empty_with_format(output_size, tensor1.options(), ACL_FORMAT_ND);
   }
++  Tensor contiguous_self = tensor1;
++  Tensor contiguous_mat2 = tensor2;
++  bool is_self_t = is_transpose_last_two_dims_v2(tensor1);
++  bool is_mat2_t = is_transpose_last_two_dims_v2(tensor2);
++
++  if(!is_self_t) {
++    contiguous_self = NpuUtils::format_contiguous(tensor1);
   }
++  if(!is_mat2_t) {
++    contiguous_mat2 = NpuUtils::format_contiguous(tensor2);
   }
++  auto func1 = [&contiguous_self]() {
       bool pass = false;
++      return std::tie(pass, contiguous_self);
   };
++  auto func2 = [&contiguous_mat2]() {
       bool pass = false;
++      return std::tie(pass, contiguous_mat2);
   };
+-  bool isSelfT = CalcuOpUtil::is_transpose_last_two_dims(self);
+-  bool isMat2T = CalcuOpUtil::is_transpose_last_two_dims(mat2);
+-
   // executing the NPU operator
   OpCommand cmd;
   cmd.Name("BatchMatMul")
       .InputWithFunc(func1)
       .InputWithFunc(func2)
       .Output(result)
++      .Attr("adj_x1", is_self_t)
++      .Attr("adj_x2", is_mat2_t)
       .Run();
   return result;
 }
++
++Tensor reshape_tensor_self(const Tensor& self, SmallVector<int64_t, SIZE>& expect_output_size) {
++  // self, expect_output: [5,6,7,17], [1,6,7,65]
++  // self permute + reshape: [5,6,7,17] -> [6,7,5,17] -> [6,7,85]
++  SmallVector<int64_t, SIZE> self_permute_idx;
++  SmallVector<int64_t, SIZE> self_batch_idx;
++
++  for (int64_t i = 0; i < self.dim(); ++i) {
++    if (i < self.dim() - 2) {
++      if (expect_output_size[i] == 1) {
++        self_batch_idx.emplace_back(i);
++        continue;
++      }
++    } else if (i == self.dim() - 1) {
++      for (int64_t j = 0; j < self_batch_idx.size(); ++j) {
++        self_permute_idx.emplace_back(self_batch_idx[j]);
++      }
++    }
++    self_permute_idx.emplace_back(i);
++  }
++  Tensor tmp_self = self.permute(self_permute_idx);
++
++  int64_t m_idx = 0;
++  SmallVector<int64_t, SIZE> tmp_self_size;
++  SmallVector<int64_t, SIZE> tmp_self_size_low;
++
++  m_idx = self.dim() - self_batch_idx.size() - 1;
++  tmp_self_size = array_to_small_vector(tmp_self.sizes());
++  tmp_self_size_low.insert(tmp_self_size_low.end(), tmp_self_size.begin(), tmp_self_size.begin() + m_idx);
++  tmp_self_size_low.emplace_back(-1);
++  tmp_self = tmp_self.reshape(tmp_self_size_low);
++  return tmp_self;
++}
++
++Tensor reshape_tensor_mat2(const Tensor& mat2, SmallVector<int64_t, SIZE>& expect_output_size) {
++  // mat2, expect_output_size: [5,6,17,65], [1,6,7,65]
++  // mat2 permute + reshape: [5,6,17,65] -> [6,5,17,65] -> [6,85,65]
++  SmallVector<int64_t, SIZE> mat2_permute_idx;
++  SmallVector<int64_t, SIZE> mat2_batch_idx;
++
++  for (int64_t i = 0; i < mat2.dim(); ++i) {
++    if (i < mat2.dim() - 2) {
++      if (expect_output_size[i] == 1) {
++        mat2_batch_idx.emplace_back(i);
++        continue;
++      }
++    } else if (i == mat2.dim() - 2) {
++      for (int64_t j = 0; j < mat2_batch_idx.size(); ++j) {
++        mat2_permute_idx.emplace_back(mat2_batch_idx[j]);
++      }
++    }
++    mat2_permute_idx.emplace_back(i);
++  }
++  Tensor tmp_mat2 = mat2.permute(mat2_permute_idx);
++
++  int64_t k_idx = 0;
++  SmallVector<int64_t, SIZE> tmp_mat2_size;
++  SmallVector<int64_t, SIZE> tmp_mat2_size_low;
++
++  k_idx = mat2.dim() - mat2_batch_idx.size() - 2;
++  tmp_mat2_size = array_to_small_vector(tmp_mat2.sizes());
++  tmp_mat2_size_low.insert(tmp_mat2_size_low.end(), tmp_mat2_size.begin(), tmp_mat2_size.begin() + k_idx);
++  tmp_mat2_size_low.insert(tmp_mat2_size_low.end(), {-1, mat2.size(-1)});
++  tmp_mat2 = tmp_mat2.reshape(tmp_mat2_size_low);
++  return tmp_mat2;
++}
++
++SmallVector<int64_t, SIZE> align_small_vector(SmallVector<int64_t, SIZE> svec,
++                                              SmallVector<int64_t, SIZE> golden_svec) {
++  // svec, golden: [6,7,65], [5,6,7,65]
++  // expect: [6,7,65] -> [1,6,7,65]
++  SmallVector<int64_t, SIZE> tmp_svec;
++  tmp_svec = svec;
++  int64_t size_to_fill = golden_svec.size() - svec.size();
++  if (size_to_fill > 0) {
++    tmp_svec.insert(tmp_svec.begin(), size_to_fill, 1);
++  }
++  return tmp_svec;
++}
++
++void expand_tensor(Tensor& self, Tensor& mat2, SmallVector<int64_t, SIZE>& expand_output_size) {
++  self = self.dim() == 1 ? self.view({1, self.size(0)}) : self;
++  mat2 = mat2.dim() == 1 ? mat2.view({mat2.size(0), 1}) : mat2;
++  int64_t m = self.size(-2);
++  int64_t k1 = self.size(-1);
++  int64_t k2 = mat2.size(-2);
++  int64_t n = mat2.size(-1);
++
++  std::vector<int64_t> expand_batch_portion(expand_output_size.begin(), expand_output_size.end() - 2);
++  std::vector<int64_t> self_expand_size(expand_batch_portion);
++  std::vector<int64_t> mat2_expand_size(expand_batch_portion);
++
++  self_expand_size.insert(self_expand_size.end(), {m, k1});
++  mat2_expand_size.insert(mat2_expand_size.end(), {k2, n});
++
++  int64_t expand_batch_product = std::accumulate(expand_batch_portion.begin(), expand_batch_portion.end(),
++                                                 1L, std::multiplies<int64_t>());
++
++  std::vector<int64_t> self_bmm_view({expand_batch_product});
++  std::vector<int64_t> mat2_bmm_view({expand_batch_product});
++  self_bmm_view.insert(self_bmm_view.end(), {m, k1});
++  mat2_bmm_view.insert(mat2_bmm_view.end(), {k2, n});
++
++  self = self.expand(self_expand_size).reshape(self_bmm_view);
++  mat2 = mat2.expand(mat2_expand_size).reshape(mat2_bmm_view);
++}
++
++Tensor bmm_v2_npu(const Tensor& self, const Tensor& mat2, IntArrayRef output_sizes) {
++  auto expect_output_size = array_to_small_vector(output_sizes);
++  auto infer_output_size = bmm_v2_output_size(self, mat2);
++  Tensor tmp_self = self;
++  Tensor tmp_mat2 = mat2;
++
++  // forward propagation
++  if (expect_output_size.empty()) {
++    // avoid some accuracy error caused by transdata
++    expand_tensor(tmp_self, tmp_mat2, infer_output_size);
++    expect_output_size = infer_output_size;
++    infer_output_size = bmm_v2_output_size(tmp_self, tmp_mat2);
++
++    auto res = pure_bmm_v2_npu(tmp_self, tmp_mat2, infer_output_size).view(expect_output_size);
++    infer_output_size = expect_output_size;
++
++    if (self.dim() == 1) {
++      // [k][b, k, n] -> [b, 1, n] -> [b, n]
++      infer_output_size.erase(infer_output_size.end() - 2);
++      return res.view(infer_output_size);
++    } else if (mat2.dim() == 1) {
++      // [b, m, k][k] -> [b, m, 1] -> [b, m]
++      infer_output_size.erase(infer_output_size.end() - 1);
++      return res.view(infer_output_size);
++    }
++    return res;
++  }
++
++  // backward propagation
++  SmallVector<int64_t, SIZE> tmp_expect_output_size = expect_output_size;
++  SmallVector<int64_t, SIZE> axis_reduce;
++  SmallVector<int64_t, SIZE> tmp_self_size;
++  SmallVector<int64_t, SIZE> tmp_mat2_size;
++
++  tmp_expect_output_size = align_small_vector(expect_output_size, infer_output_size);
++  for (int i = 0; i < tmp_expect_output_size.size(); ++i) {
++    if (tmp_expect_output_size[i] != infer_output_size[i]) {
++      axis_reduce.emplace_back(i);
++    }
++  }
++
++  // no reduce_sum
++  if (axis_reduce.empty()) {
++    // avoid some accuracy error caused by transdata
++    expand_tensor(tmp_self, tmp_mat2, infer_output_size);
++    infer_output_size = bmm_v2_output_size(tmp_self, tmp_mat2);
++    return pure_bmm_v2_npu(tmp_self, tmp_mat2, infer_output_size).view(expect_output_size);
++  }
++
++  // reduce sum without accuracy error
++  tmp_self_size = align_small_vector(array_to_small_vector(self.sizes()), infer_output_size);
++  tmp_mat2_size = align_small_vector(array_to_small_vector(mat2.sizes()), infer_output_size);
++  tmp_self = self.reshape(tmp_self_size);
++  tmp_mat2 = mat2.reshape(tmp_mat2_size);
++  tmp_self = reshape_tensor_self(tmp_self, tmp_expect_output_size);
++  tmp_mat2 = reshape_tensor_mat2(tmp_mat2, tmp_expect_output_size);
++  infer_output_size = bmm_v2_output_size(tmp_self, tmp_mat2);
++  // avoid some accuracy error caused by transdata
++  expand_tensor(tmp_self, tmp_mat2, infer_output_size);
++  infer_output_size = bmm_v2_output_size(tmp_self, tmp_mat2);
++  return pure_bmm_v2_npu(tmp_self, tmp_mat2, infer_output_size).view(expect_output_size);
++}
++
 } // namespace native
 } // namespace at

Msrc/aten/src/ATen/native/npu/BoundingBoxDecodeKernelNpu.cpp+3-6

	@@ -14,9 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include <ATen/ATen.h>	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19	-#include "ATen/native/npu/utils/OpTemplate.h"
20		19
21	namespace at {	20	namespace at {
22	namespace native {	21	namespace native {
	@@ -59,9 +58,7 @@ Tensor bounding_box_decode_npu(
59	double wh_ratio_clip) {	58	double wh_ratio_clip) {
60	SmallVector<int64_t, SIZE> outputSize = {rois.size(0), 4};	59	SmallVector<int64_t, SIZE> outputSize = {rois.size(0), 4};
61	// construct the output tensor of the NPU	60	// construct the output tensor of the NPU
62	- Tensor result = at::empty_with_format(	61	+ Tensor result = OpPreparation::ApplyTensor(rois, outputSize);
63	- outputSize, rois.options(), CalcuOpUtil::get_tensor_npu_format(rois));
64	-
65	SmallVector<float, SIZE> means = {	62	SmallVector<float, SIZE> means = {
66	static_cast<float>(means0),	63	static_cast<float>(means0),
67	static_cast<float>(means1),	64	static_cast<float>(means1),

Msrc/aten/src/ATen/native/npu/BoundingBoxEncodeKernelNpu.cpp+1-5

	@@ -51,11 +51,7 @@ Tensor bounding_box_encode_npu(
51	double stds2,	51	double stds2,
52	double stds3) {	52	double stds3) {
53	// construct the output tensor of the NPU	53	// construct the output tensor of the NPU
54	- Tensor delats = at::empty_with_format(	54	+ Tensor delats = OpPreparation::ApplyTensor(anchor_box, {anchor_box.size(0), 4});
55	- {anchor_box.size(0), 4},
56	- anchor_box.options(),
57	- CalcuOpUtil::get_tensor_npu_format(anchor_box));
58	-
59	SmallVector<float, SIZE> means = {	55	SmallVector<float, SIZE> means = {
60	static_cast<float>(means0),	56	static_cast<float>(means0),
61	static_cast<float>(means1),	57	static_cast<float>(means1),

Msrc/aten/src/ATen/native/npu/BroadcastKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	#include "c10/npu/OptionsManager.h"	19	#include "c10/npu/OptionsManager.h"
19		20
20	namespace at {	21	namespace at {

Msrc/aten/src/ATen/native/npu/CastKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/CatKernelNpu.cpp+1-0

	@@ -16,6 +16,7 @@
16		16
17	#include "c10/npu/OptionsManager.h"	17	#include "c10/npu/OptionsManager.h"
18	#include "ATen/native/npu/utils/OpAdapter.h"	18	#include "ATen/native/npu/utils/OpAdapter.h"
		19	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19		20
20	namespace at {	21	namespace at {
21	namespace native {	22	namespace native {

Msrc/aten/src/ATen/native/npu/CdistBackwardKernelNpu.cpp+2-10

	@@ -12,10 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"
17	-#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	-#include "ATen/native/npu/utils/NpuUtils.h"
19		16
20	namespace at {	17	namespace at {
21	namespace native {	18	namespace native {
	@@ -81,12 +78,7 @@ Tensor _cdist_backward_npu(
81		78
82	//Executing the NPU operator.	79	//Executing the NPU operator.
83	auto outputSize = input_same_output_size(x1);	80	auto outputSize = input_same_output_size(x1);
84	-	81	+ Tensor result = OpPreparation::ApplyTensor(tensor1_broadcast, outputSize);
85	- Tensor result = at::empty_with_format(
86	- outputSize,
87	- tensor1_broadcast.options(),
88	- CalcuOpUtil::get_tensor_npu_format(tensor1_broadcast));
89	-
90	OpCommand cmd;	82	OpCommand cmd;
91	cmd.Name("CdistGrad")	83	cmd.Name("CdistGrad")
92	.Input(grad_broadcast)	84	.Input(grad_broadcast)

Msrc/aten/src/ATen/native/npu/CeilKernelNpu.cpp+2-13

	@@ -33,24 +33,13 @@ Tensor& ceil_out_npu(Tensor& result, const Tensor& self) {
33	}	33	}
34		34
35	Tensor ceil_npu(const Tensor& self) {	35	Tensor ceil_npu(const Tensor& self) {
36	- // calculate the output size	36	+ Tensor result = OpPreparation::ApplyTensor(self);
37	- auto outputSize = input_same_output_size(self);
38	-
39	- // construct the output tensor of the NPU
40	- Tensor result = at::empty_with_format(
41	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
42	-
43	- // calculate the output result of the NPU
44	ceil_out_npu(result, self);	37	ceil_out_npu(result, self);
45	-
46	return result;	38	return result;
47	}	39	}
48		40
49	Tensor& ceil_npu_(Tensor& self) {	41	Tensor& ceil_npu_(Tensor& self) {
50	- SmallVector<Tensor, N> inputs = {self};	42	+ OpPreparation::CheckMemory({self}, {self});
51	- SmallVector<Tensor, N> outputs = {self};
52	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
53	-
54	if (!NpuUtils::check_match(&self)) {	43	if (!NpuUtils::check_match(&self)) {
55	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	44	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
56	Tensor result = ceil_out_npu(contiguousSelf, contiguousSelf);	45	Tensor result = ceil_out_npu(contiguousSelf, contiguousSelf);

Rsrc/aten/src/ATen/native/npu/TraceKernelNpu.cpp→src/aten/src/ATen/native/npu/CeluKernelNpu.cpp+36-27

	@@ -1,4 +1,6 @@
1	-// Copyright (c) 2020, Huawei Technologies.All rights reserved.	1	+// Copyright (c) 2020 Huawei Technologies Co., Ltd
		2	+// Copyright (c) 2019, Facebook CORPORATION.
		3	+// All rights reserved.
2	//	4	//
3	// Licensed under the BSD 3-Clause License (the "License");	5	// Licensed under the BSD 3-Clause License (the "License");
4	// you may not use this file except in compliance with the License.	6	// you may not use this file except in compliance with the License.
	@@ -12,39 +14,46 @@
12	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
13	// limitations under the License.	15	// limitations under the License.
14		16
15	-#include "ATen/native/npu/utils/CalcuOpUtil.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
17	-#include "ATen/native/npu/utils/NpuUtils.h"
18		18
19	namespace at {	19	namespace at {
20	namespace native {	20	namespace native {
21	using namespace at::native::npu;	21	using namespace at::native::npu;
22	-SmallVector<NPUTensorDesc, N> trace_npu_input(const SmallVector<Tensor, N>& inputTensor) {
23	- return CalcuOpUtil::create_npu_input_tensor_desc(inputTensor);
24	-}
25		22
26	-SmallVector<NPUTensorDesc, N> trace_npu_output(const SmallVector<Tensor, N>& outputTensor) {	23	+Tensor celu_out_npu_nocheck(Tensor& result, const Tensor& self, Scalar alpha) {
27	- return CalcuOpUtil::create_npu_output_tensor_desc(outputTensor);	24	+ float alpha3 = 1.0;
28	-}	25	+ OpCommand cmd;
29	-	26	+ cmd.Name("Celu")
30	-SmallVector<NPUAttrDesc, N> trace_npu_attr() {	27	+ .Input(self)
31	- SmallVector<NPUAttrDesc, N> attrs = {};	28	+ .Output(result)
32	- return attrs;	29	+ .Attr("alpha1", alpha)
33	-}	30	+ .Attr("alpha2", alpha)
34	-	31	+ .Attr("alpha3", alpha3)
35	-Tensor& trace_out_npu(Tensor& result, const Tensor& self) {	32	+ .Run();
36	- auto inputs = trace_npu_input({self});
37	- auto outputs = trace_npu_output({result});
38	- auto attrs = trace_npu_attr();
39	- CalcuOpUtil::execute_npu_operate("Trace", inputs, outputs, attrs);
40	return result;	33	return result;
41	}	34	}
42		35
43	-Tensor trace_npu(const Tensor& self) {	36	+Tensor celu_out_npu(Tensor& result, const Tensor& self, Scalar alpha) {
44	- auto outputSize = trace_npu_output_size(self);	37	+ OpPipeWithDefinedOut pipe;
45	- Tensor result = at::empty_with_format(outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));	38	+ return pipe.CheckMemory({self}, {result})
46	- trace_out_npu(result, self);	39	+ .Func([&self, &alpha](Tensor& result){celu_out_npu_nocheck(result, self, alpha);})
47	- return result.reshape({});	40	+ .Call(result);
48	}	41	}
		42	+
		43	+Tensor celu_npu(const Tensor& self, Scalar alpha) {
		44	+ // construct the output tensor of the NPU
		45	+ Tensor result = OpPreparation::ApplyTensor(self);
		46	+
		47	+ // calculate the output result of the NPU
		48	+ celu_out_npu(result, self, alpha);
		49	+
		50	+ return result;
49	}	51	}
50	-}	52	+
		53	+Tensor& celu_npu_(Tensor& self, Scalar alpha) {
		54	+ celu_out_npu(self, self, alpha);
		55	+ return self;
		56	+}
		57	+
		58	+} // namespace native
		59	+} // namespace at

Msrc/aten/src/ATen/native/npu/CholeskyKernelNpu.cpp+0-1

	@@ -14,7 +14,6 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
19	#include "ATen/native/npu/utils/NpuUtils.h"	18	#include "ATen/native/npu/utils/NpuUtils.h"
20	#include "ATen/native/npu/utils/OpTemplate.h"	19	#include "ATen/native/npu/utils/OpTemplate.h"

Msrc/aten/src/ATen/native/npu/ClampKernelNpu.cpp+14-56

	@@ -16,8 +16,6 @@
16		16
17	#include <climits>	17	#include <climits>
18	#include <float.h>	18	#include <float.h>
19	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
20	-#include "ATen/native/npu/utils/OpTemplate.h"
21	#include "ATen/native/npu/utils/OpAdapter.h"	19	#include "ATen/native/npu/utils/OpAdapter.h"
22		20
23	namespace at {	21	namespace at {
	@@ -40,14 +38,11 @@ Tensor& clamp_min_out_npu_nocheck(
40	max = NPU_HALF_MAX;	38	max = NPU_HALF_MAX;
41	}	39	}
42		40
43	- Tensor minTensor = CalcuOpUtil::CopyScalarToDevice(min, self.scalar_type());
44	- Tensor maxTensor = CalcuOpUtil::CopyScalarToDevice(max, self.scalar_type());
45	-
46	OpCommand cmd;	41	OpCommand cmd;
47	cmd.Name("ClipByValue")	42	cmd.Name("ClipByValue")
48	.Input(self)	43	.Input(self)
49	- .Input(minTensor)	44	+ .Input(min, self.scalar_type())
50	- .Input(maxTensor)	45	+ .Input(max, self.scalar_type())
51	.Output(result)	46	.Output(result)
52	.Run();	47	.Run();
53	return result;	48	return result;
	@@ -60,9 +55,7 @@ Tensor& clamp_min_out_npu(
60	OpPreparation::CheckOut(	55	OpPreparation::CheckOut(
61	{self},	56	{self},
62	result,	57	result,
63	- CalcuOpUtil::get_tensor_npu_format(self),	58	+ self);
64	- self.scalar_type(),
65	- self.sizes());
66		59
67	OpPipeWithDefinedOut pipe;	60	OpPipeWithDefinedOut pipe;
68	return pipe.CheckMemory({self}, {result})	61	return pipe.CheckMemory({self}, {result})
	@@ -84,15 +77,12 @@ Tensor& clamp_max_out_npu(Tensor& result, const Tensor& self, Scalar max) {
84	} else {	77	} else {
85	min = NPU_HALF_MIN;	78	min = NPU_HALF_MIN;
86	}	79	}
87	-	80	+
88	- Tensor minTensor = CalcuOpUtil::CopyScalarToDevice(min, self.scalar_type());
89	- Tensor maxTensor = CalcuOpUtil::CopyScalarToDevice(max, self.scalar_type());
90	-
91	OpCommand cmd;	81	OpCommand cmd;
92	cmd.Name("ClipByValue")	82	cmd.Name("ClipByValue")
93	.Input(self)	83	.Input(self)
94	- .Input(minTensor)	84	+ .Input(min, self.scalar_type())
95	- .Input(maxTensor)	85	+ .Input(max, self.scalar_type())
96	.Output(result)	86	.Output(result)
97	.Run();	87	.Run();
98	return result;	88	return result;
	@@ -111,15 +101,12 @@ Tensor& clamp_out_npu_nocheck(
111	Scalar minScalar = min.value();	101	Scalar minScalar = min.value();
112	clamp_min_out_npu(result, self, minScalar);	102	clamp_min_out_npu(result, self, minScalar);
113		103
114	- } else {	104	+ } else {
115	- Tensor minTensor = CalcuOpUtil::CopyScalarToDevice(min.value(), self.scalar_type());
116	- Tensor maxTensor = CalcuOpUtil::CopyScalarToDevice(max.value(), self.scalar_type());
117	-
118	OpCommand cmd;	105	OpCommand cmd;
119	cmd.Name("ClipByValue")	106	cmd.Name("ClipByValue")
120	.Input(self)	107	.Input(self)
121	- .Input(minTensor)	108	+ .Input(min.value(), self.scalar_type())
122	- .Input(maxTensor)	109	+ .Input(max.value(), self.scalar_type())
123	.Output(result)	110	.Output(result)
124	.Run();	111	.Run();
125	}	112	}
	@@ -135,9 +122,7 @@ Tensor& clamp_out_npu(
135	OpPreparation::CheckOut(	122	OpPreparation::CheckOut(
136	{self},	123	{self},
137	result,	124	result,
138	- CalcuOpUtil::get_tensor_npu_format(self),	125	+ self);
139	- self.scalar_type(),
140	- self.sizes());
141		126
142	OpPipeWithDefinedOut pipe;	127	OpPipeWithDefinedOut pipe;
143	return pipe.CheckMemory({self}, {result})	128	return pipe.CheckMemory({self}, {result})
	@@ -148,16 +133,8 @@ Tensor& clamp_out_npu(
148	}	133	}
149		134
150	Tensor clamp_min_npu(const Tensor& self, Scalar min) {	135	Tensor clamp_min_npu(const Tensor& self, Scalar min) {
151	- // calculate the output size	136	+ Tensor result = OpPreparation::ApplyTensor(self);
152	- auto outputSize = input_same_output_size(self);
153	-
154	- // construct the output tensor of the NPU
155	- Tensor result = at::empty_with_format(
156	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
157	-
158	- // calculate the output result of the NPU
159	clamp_min_out_npu_nocheck(result, self, min);	137	clamp_min_out_npu_nocheck(result, self, min);
160	-
161	return result;	138	return result;
162	}	139	}
163		140
	@@ -168,24 +145,14 @@ Tensor& clamp_min_npu_(Tensor& self, Scalar min) {
168	}	145	}
169		146
170	Tensor clamp_max_npu(const Tensor& self, Scalar max) {	147	Tensor clamp_max_npu(const Tensor& self, Scalar max) {
171	- // calculate the output size	148	+ Tensor result = OpPreparation::ApplyTensor(self);
172	- auto outputSize = input_same_output_size(self);
173	-
174	- // construct the output tensor of the NPU
175	- Tensor result = at::empty_with_format(
176	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
177	-
178	- // calculate the output result of the NPU
179	clamp_max_out_npu(result, self, max);	149	clamp_max_out_npu(result, self, max);
180		150
181	return result;	151	return result;
182	}	152	}
183		153
184	Tensor& clamp_max_npu_(Tensor& self, Scalar max) {	154	Tensor& clamp_max_npu_(Tensor& self, Scalar max) {
185	- SmallVector<Tensor, N> inputs = {self};	155	+ OpPreparation::CheckMemory({self}, {self});
186	- SmallVector<Tensor, N> outputs = {self};
187	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
188	-
189	if (!NpuUtils::check_match(&self)) {	156	if (!NpuUtils::check_match(&self)) {
190	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	157	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
191	Tensor result = clamp_max_out_npu(contiguousSelf, contiguousSelf, max);	158	Tensor result = clamp_max_out_npu(contiguousSelf, contiguousSelf, max);
	@@ -201,22 +168,13 @@ Tensor clamp_npu(
201	const Tensor& self,	168	const Tensor& self,
202	optional<Scalar> min,	169	optional<Scalar> min,
203	optional<Scalar> max) {	170	optional<Scalar> max) {
204	- // calculate the output size	171	+ Tensor result = OpPreparation::ApplyTensor(self);
205	- auto outputSize = input_same_output_size(self);
206	-
207	- // construct the output tensor of the NPU
208	- Tensor result = at::empty_with_format(
209	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
210	-
211	- // calculate the output result of the NPU
212	clamp_out_npu_nocheck(result, self, min, max);	172	clamp_out_npu_nocheck(result, self, min, max);
213	-
214	return result;	173	return result;
215	}	174	}
216		175
217	Tensor& clamp_npu_(Tensor& self, optional<Scalar> min, optional<Scalar> max) {	176	Tensor& clamp_npu_(Tensor& self, optional<Scalar> min, optional<Scalar> max) {
218	clamp_out_npu(self, self, min, max);	177	clamp_out_npu(self, self, min, max);
219	-
220	return self;	178	return self;
221	}	179	}
222		180

Msrc/aten/src/ATen/native/npu/ConfusionTransposeKernelNpu.cpp+2-4

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -36,8 +35,7 @@ Tensor confusion_transpose_npu(
36	}	35	}
37		36
38	// construct the output tensor of the NPU	37	// construct the output tensor of the NPU
39	- Tensor result = at::empty_with_format(output_size, self.options(), CalcuOpUtil::get_tensor_npu_format(self));	38	+ Tensor result = OpPreparation::ApplyTensor(self, output_size);
40	-
41	OpCommand cmd;	39	OpCommand cmd;
42	cmd.Name("ConfusionTransposeD")	40	cmd.Name("ConfusionTransposeD")
43	.Input(self)	41	.Input(self)

Msrc/aten/src/ATen/native/npu/ConvTbcKernelNpu.cpp+24-66

 // See the License for the specific language governing permissions and
 // limitations under the License.
++#include "ATen/native/npu/utils/OpAdapter.h"
+-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
+-#include "ATen/native/npu/utils/NpuUtils.h"
 namespace at {
 namespace native {
 using namespace at::native::npu;
+-SmallVector<NPUTensorDesc, N> conv_tbc_npu_input(
+-    const SmallVector<Tensor, N>& inputTensor) {
+-  SmallVector<Tensor, N> inputTensors;
+-  for (int i = 0; i < inputTensor.size(); i++) {
+-    if (inputTensor[i].defined()) {
+-      inputTensors.emplace_back(inputTensor[i]);
+-    }
+-  }
+-
+-  return CalcuOpUtil::create_npu_input_tensor_desc(inputTensors);
+-}
+-
+-SmallVector<NPUTensorDesc, N> conv_tbc_npu_output(
+-    const SmallVector<Tensor, N>& outputTensor) {
+-  return CalcuOpUtil::create_npu_output_tensor_desc(outputTensor);
+-}
+-
+-SmallVector<NPUAttrDesc, N> conv_tbc_npu_attr(int64_t pad) {
+-  SmallVector<int64_t, N> paddings = {0, 0, pad, pad};
+-  SmallVector<int64_t, N> stridesSize = {1, 1, 1, 1};
+-  SmallVector<int64_t, N> dilations = {1, 1, 1, 1};
+-
+-  string dataFormat = "NCHW";
+-
+-  NPUAttrDesc npuAttrPads = NPUAttrDesc("pads", paddings);
+-  NPUAttrDesc npuAttrStrides = NPUAttrDesc("strides", stridesSize);
+-  NPUAttrDesc npuAttrDilations = NPUAttrDesc("dilations", dilations);
+-  NPUAttrDesc npuAttrDataFormat = NPUAttrDesc("data_format", dataFormat);
+-
+-  SmallVector<NPUAttrDesc, N> attrs = {
+-      npuAttrPads, npuAttrStrides, npuAttrDilations, npuAttrDataFormat};
+-
+-  return attrs;
+-}
+-
+-Tensor& conv_tbc_out_npu(
+-    Tensor& result,
+-    const Tensor& self,
+-    const Tensor& weight,
+-    const Tensor& bias,
+-    int64_t pad) {
+-  // constructs the input and output NPUTensorDesc
+-
+-  auto inputs = conv_tbc_npu_input(
+-      {self.transpose(0, 2).transpose(0, 1).unsqueeze(2),
+-       weight.transpose(0, 2).unsqueeze(2),
+-       bias});
+-
+-  auto outputs = conv_tbc_npu_output({result});
+-
+-  // constructs the attr of the NPUAttrDesc
+-  auto attrs = conv_tbc_npu_attr(pad);
+-  // executing the NPU operator
+-  CalcuOpUtil::execute_npu_operate("Conv2D", inputs, outputs, attrs);
+-
+-  return result;
+-}
+-
 Tensor conv_tbc_npu(
     const Tensor& self,
     const Tensor& weight,
       "the weight tensor (output channels).");
   // calculate the output size
++  int64_t Co = weight.size(2);
++  int64_t Wo = (self.size(0) + 2 * pad - (weight.size(0) - 1) - 1) + 1;
++
++  SmallVector<int64_t, SIZE> outputSize = {self.size(1), Co, 1, Wo};
   // construct the output tensor of the NPU
++  Tensor result = OpPreparation::ApplyTensorWithFormat(self, outputSize, ACL_FORMAT_NCHW);
+-      at::empty_with_format(outputSize, self.options(), ACL_FORMAT_NCHW);
++  SmallVector<int64_t, N> paddings = {0, 0, pad, pad};
++  SmallVector<int64_t, N> stridesSize = {1, 1, 1, 1};
++  SmallVector<int64_t, N> dilations = {1, 1, 1, 1};
++
++  Tensor self_tensor = self.transpose(0, 2).transpose(0, 1).unsqueeze(2);
++  Tensor weight_tensor = weight.transpose(0, 2).unsqueeze(2);
++
++  OpCommand cmd;
++  cmd.Name("Conv2D")
++    .Input(self_tensor)
++    .Input(weight_tensor)
++    .Input(bias)
++    .Output(result)
++    .Attr("pads", paddings)
++    .Attr("strides", stridesSize)
++    .Attr("dilations", dilations)
++    .Attr("data_format", (string)"NCHW")
++    .Run();
   result = result.squeeze(2).transpose(0, 2).transpose(1, 2);
   return result;

Msrc/aten/src/ATen/native/npu/CosKernelNpu.cpp+3-17

	@@ -12,8 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"
17		16
18	namespace at {	17	namespace at {
19	namespace native {	18	namespace native {
	@@ -32,26 +31,13 @@ Tensor& cos_out_npu(
32	}	31	}
33		32
34	Tensor cos_npu(const Tensor& self) {	33	Tensor cos_npu(const Tensor& self) {
35	- // calculate the output size	34	+ Tensor result = OpPreparation::ApplyTensor(self);
36	- auto outputSize = input_same_output_size(self);
37	-
38	- // construct the output tensor of the NPU
39	- Tensor result = at::empty_with_format(
40	- outputSize,
41	- self.options(),
42	- CalcuOpUtil::get_tensor_npu_format(self));
43	-
44	- // calculate the output result of the NPU
45	cos_out_npu(result, self);	35	cos_out_npu(result, self);
46	-
47	return result;	36	return result;
48	}	37	}
49		38
50	Tensor& cos_npu_(Tensor& self) {	39	Tensor& cos_npu_(Tensor& self) {
51	- SmallVector<Tensor, N> inputs = {self};	40	+ OpPreparation::CheckMemory({self}, {self});
52	- SmallVector<Tensor, N> outputs = {self};
53	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
54	-
55	if (!NpuUtils::check_match(&self)) {	41	if (!NpuUtils::check_match(&self)) {
56	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	42	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
57	Tensor result = cos_out_npu(contiguousSelf, contiguousSelf);	43	Tensor result = cos_out_npu(contiguousSelf, contiguousSelf);

Msrc/aten/src/ATen/native/npu/CoshKernelNpu.cpp+3-12

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -32,9 +31,7 @@ Tensor& cosh_out_npu(Tensor& result, const Tensor& self) {
32	}	31	}
33		32
34	Tensor& cosh_npu_(Tensor& self) {	33	Tensor& cosh_npu_(Tensor& self) {
35	- SmallVector<Tensor, N> inputs = {self};	34	+ OpPreparation::CheckMemory({self}, {self});
36	- SmallVector<Tensor, N> outputs = {self};
37	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
38		35
39	if (!NpuUtils::check_match(&self)) {	36	if (!NpuUtils::check_match(&self)) {
40	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	37	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
	@@ -48,13 +45,7 @@ Tensor& cosh_npu_(Tensor& self) {
48	}	45	}
49		46
50	Tensor cosh_npu(const Tensor& self) {	47	Tensor cosh_npu(const Tensor& self) {
51	- // calculate the output size	48	+ Tensor result = OpPreparation::ApplyTensor(self);
52	- auto outputSize = input_same_output_size(self);
53	-
54	- // construct the output tensor of the NPU
55	- Tensor result = at::empty_with_format(
56	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
57	-
58	// calculate the output result of the NPU	49	// calculate the output result of the NPU
59	cosh_out_npu(result, self);	50	cosh_out_npu(result, self);
60	return result;	51	return result;

Msrc/aten/src/ATen/native/npu/CrossKernelNpu.cpp+2-2

	@@ -12,8 +12,8 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"	16	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
17		17
18	namespace at {	18	namespace at {
19	namespace native {	19	namespace native {

Msrc/aten/src/ATen/native/npu/CtcLossBackwardKernelNpu.cpp+18-19

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -50,30 +49,29 @@ Tensor ctc_loss_backward_npu(
50	if (logAlpha.scalar_type() == ScalarType::Half) {	49	if (logAlpha.scalar_type() == ScalarType::Half) {
51	logAlphaNeed = logAlphaNeed.to(ScalarType::Float);	50	logAlphaNeed = logAlphaNeed.to(ScalarType::Float);
52	}	51	}
		52	+
		53	+ Tensor targetsCast = targets;
		54	+ if(targets.scalar_type() == ScalarType::Long){
		55	+ targetsCast = targetsCast.to(ScalarType::Int);
		56	+ }
		57	+
		58	+ auto inputLengthsTensor = at::tensor(inputLengths, targetsCast.options().dtype(at::kInt));
		59	+ auto targetLengthsTensor = at::tensor(targetLengths, targetsCast.options().dtype(at::kInt));
53		60
54	- // IntArrayRef to Tensor	61	+ auto outputSize = {logProbs.size(1), logProbs.size(0), logProbs.size(2)};
55	- auto inputLengthsTensor = at::tensor(inputLengths, targets.options().dtype(at::kLong));
56	- auto targetLengthsTensor = at::tensor(targetLengths, targets.options().dtype(at::kLong));
57	-
58	- // calculate the output size
59	- auto outputSize = input_same_output_size(logProbs);
60		62
61	// construct the output tensor of the NPU	63	// construct the output tensor of the NPU
62	- Tensor grad = at::empty_with_format(	64	+ Tensor grad = OpPreparation::ApplyTensor(logProbsNeed, outputSize);
63	- outputSize,
64	- logProbsNeed.options(),
65	- CalcuOpUtil::get_tensor_npu_format(logProbsNeed));
66	-
67	// calculate the output result of the NPU	65	// calculate the output result of the NPU
68	OpCommand cmd;	66	OpCommand cmd;
69	cmd.Name("CTCLossV2Grad")	67	cmd.Name("CTCLossV2Grad")
70	.Input(gradOutNeed)	68	.Input(gradOutNeed)
71	.Input(logProbsNeed)	69	.Input(logProbsNeed)
72	- .Input(targets)	70	+ .Input(targetsCast)
73	- .Input(negLogLikelihoodNeed)
74	- .Input(logAlphaNeed)
75	.Input(inputLengthsTensor)	71	.Input(inputLengthsTensor)
76	- .Input(targetLengthsTensor)	72	+ .Input(targetLengthsTensor)
		73	+ .Input(negLogLikelihoodNeed)
		74	+ .Input(logAlphaNeed)
77	.Output(grad)	75	.Output(grad)
78	.Attr("blank", blank)	76	.Attr("blank", blank)
79	.Attr("zero_infinity", zeroInfinity)	77	.Attr("zero_infinity", zeroInfinity)
	@@ -82,8 +80,9 @@ Tensor ctc_loss_backward_npu(
82	if (gradOut.scalar_type() == ScalarType::Half) {	80	if (gradOut.scalar_type() == ScalarType::Half) {
83	grad = grad.to(ScalarType::Half);	81	grad = grad.to(ScalarType::Half);
84	}	82	}
85	-	83	+
86	- return grad;	84	+ //return grad;
		85	+ return grad.permute({1,0,2});
87	}	86	}
88	} // namespace native	87	} // namespace native
89	} // namespace at	88	} // namespace at

Msrc/aten/src/ATen/native/npu/CtcLossKernelNpu.cpp+12-6

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19		19
20	namespace at {	20	namespace at {
21	namespace native {	21	namespace native {
	@@ -33,12 +33,18 @@ std::tuple<Tensor, Tensor> ctc_loss_npu(
33	logProbsNeed = logProbsNeed.to(ScalarType::Float);	33	logProbsNeed = logProbsNeed.to(ScalarType::Float);
34	}	34	}
35		35
		36	+ //Aicore supports only the int type
		37	+ Tensor targetsCast = targets;
		38	+ if(targets.scalar_type() == ScalarType::Long){
		39	+ targetsCast = targetsCast.to(ScalarType::Int);
		40	+ }
		41	+
36	// IntArrayRef to Tensor	42	// IntArrayRef to Tensor
37	- auto inputLengthsTensor = at::tensor(inputLengths, targets.options().dtype(at::kLong));	43	+ auto inputLengthsTensor = at::tensor(inputLengths, targetsCast.options());
38	- auto targetLengthsTensor = at::tensor(targetLengths, targets.options().dtype(at::kLong));	44	+ auto targetLengthsTensor = at::tensor(targetLengths, targetsCast.options());
39		45
40	// calculate the output size	46	// calculate the output size
41	- auto outputSizes = ctc_loss_npu_output_size(logProbs, targets, targetLengths);	47	+ auto outputSizes = ctc_loss_npu_output_size(logProbs, targetsCast, targetLengths);
42		48
43	// construct the output tensor of the NPU	49	// construct the output tensor of the NPU
44	Tensor negLogLikelihood = at::empty_with_format(	50	Tensor negLogLikelihood = at::empty_with_format(
	@@ -55,7 +61,7 @@ std::tuple<Tensor, Tensor> ctc_loss_npu(
55	OpCommand cmd;	61	OpCommand cmd;
56	cmd.Name("CTCLossV2")	62	cmd.Name("CTCLossV2")
57	.Input(logProbsNeed)	63	.Input(logProbsNeed)
58	- .Input(targets)	64	+ .Input(targetsCast)
59	.Input(inputLengthsTensor)	65	.Input(inputLengthsTensor)
60	.Input(targetLengthsTensor)	66	.Input(targetLengthsTensor)
61	.Output(negLogLikelihood)	67	.Output(negLogLikelihood)

Msrc/aten/src/ATen/native/npu/CumprodKernelNpu.cpp+1-0

	@@ -13,6 +13,7 @@
13	// limitations under the License.	13	// limitations under the License.
14		14
15	#include "ATen/native/npu/utils/OpAdapter.h"	15	#include "ATen/native/npu/utils/OpAdapter.h"
		16	+#include<ATen/NamedTensorUtils.h>
16		17
17	namespace at {	18	namespace at {
18	namespace native {	19	namespace native {

Msrc/aten/src/ATen/native/npu/DiagKernelNpu.cpp+1-2

	@@ -73,8 +73,7 @@ Tensor& diag_out_npu(Tensor& result, const Tensor& self, int64_t diagonal) {
73	OpPreparation::CheckOut(	73	OpPreparation::CheckOut(
74	{self},	74	{self},
75	result,	75	result,
76	- CalcuOpUtil::get_tensor_npu_format(self),	76	+ self,
77	- self.scalar_type(),
78	outputSize);	77	outputSize);
79		78
80	OpPipeWithDefinedOut pipe;	79	OpPipeWithDefinedOut pipe;

Msrc/aten/src/ATen/native/npu/DivKernelNpu.cpp+2-2

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19		19
20	namespace at {	20	namespace at {
21	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/DotKernelNpu.cpp+2-6

	@@ -12,10 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/CalcuOpUtil.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
17	-#include "ATen/native/npu/utils/NpuUtils.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		16
20	namespace at {	17	namespace at {
21	namespace native {	18	namespace native {
	@@ -34,9 +31,8 @@ Tensor& dot_out_npu(Tensor& result, const Tensor& self, const Tensor& tensor) {
34	return result;	31	return result;
35	}	32	}
36	Tensor dot_npu(const Tensor& self, const Tensor& tensor) {	33	Tensor dot_npu(const Tensor& self, const Tensor& tensor) {
37	- // calculate the output size
38	SmallVector<int64_t, SIZE> outputSize = dot_npu_output_size(self, tensor);	34	SmallVector<int64_t, SIZE> outputSize = dot_npu_output_size(self, tensor);
39	- Tensor result = at::empty_with_format(outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));	35	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
40	dot_out_npu(result, self, tensor);	36	dot_out_npu(result, self, tensor);
41	return result;	37	return result;
42	}	38	}

Msrc/aten/src/ATen/native/npu/DropoutBackwardKernelNpu.cpp+2-7

	@@ -13,9 +13,8 @@
13	// See the License for the specific language governing permissions and	13	// See the License for the specific language governing permissions and
14	// limitations under the License.	14	// limitations under the License.
15		15
		16	+#include "ATen/native/npu/utils/OpAdapter.h"
16	#include "ATen/native/npu/utils/CalcuOpUtil.h"	17	#include "ATen/native/npu/utils/CalcuOpUtil.h"
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -31,12 +30,8 @@ Tensor dropout_backward_npu(
31	TORCH_CHECK(	30	TORCH_CHECK(
32	mask.scalar_type() == at::ScalarType::Byte,	31	mask.scalar_type() == at::ScalarType::Byte,
33	"mask should be torch.uint8 dtype");	32	"mask should be torch.uint8 dtype");
34	- auto outputSize = input_same_output_size(grad_output);
35	double retain = 1. - scale;	33	double retain = 1. - scale;
36	- Tensor result = at::empty_with_format(	34	+ Tensor result = OpPreparation::ApplyTensor(grad_output);
37	- outputSize,
38	- grad_output.options(),
39	- CalcuOpUtil::get_tensor_npu_format(grad_output));
40	Tensor prob =	35	Tensor prob =
41	CalcuOpUtil::CopyScalarToDevice(retain, grad_output.scalar_type());	36	CalcuOpUtil::CopyScalarToDevice(retain, grad_output.scalar_type());
42		37

Msrc/aten/src/ATen/native/npu/DropoutV2BackwardKernelNpu.cpp+2-5

	@@ -13,8 +13,7 @@
13	// See the License for the specific language governing permissions and	13	// See the License for the specific language governing permissions and
14	// limitations under the License.	14	// limitations under the License.
15		15
16	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	16	+#include "ATen/native/npu/utils/OpAdapter.h"
17	-#include "ATen/native/npu/utils/OpTemplate.h"
18		17
19	namespace at {	18	namespace at {
20	namespace native {	19	namespace native {
	@@ -53,9 +52,7 @@ Tensor dropout_v2_backward_npu(const Tensor& grad_output, const Tensor& mask, do
53	if (maskCopy.scalar_type() == ScalarType::Byte){	52	if (maskCopy.scalar_type() == ScalarType::Byte){
54	maskCopy = maskCopy.to(ScalarType::Half);	53	maskCopy = maskCopy.to(ScalarType::Half);
55	}	54	}
56	- auto outputSize = input_same_output_size(grad_output);	55	+ auto result = OpPreparation::ApplyTensor(grad_output);
57	- auto result = at::empty_with_format(
58	- outputSize, grad_output.options(), CalcuOpUtil::get_tensor_npu_format(grad_output));
59	dropout_v2_backward_out_npu(result, grad_output, maskCopy, p);	56	dropout_v2_backward_out_npu(result, grad_output, maskCopy, p);
60		57
61	return result;	58	return result;

Msrc/aten/src/ATen/native/npu/DropoutV2KernelNpu.cpp+3-7

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -46,11 +45,8 @@ tuple <Tensor, Tensor, Tensor> dropout_v2_npu(const Tensor& self, Tensor& seed,
46	Tensor formatCastOfSelf = OpPreparation::CastBackToOriFormat(self);	45	Tensor formatCastOfSelf = OpPreparation::CastBackToOriFormat(self);
47	Tensor formatCastOfSeed = OpPreparation::CastBackToOriFormat(seed);	46	Tensor formatCastOfSeed = OpPreparation::CastBackToOriFormat(seed);
48		47
49	- Tensor result = at::empty_with_format(	48	+ Tensor result = OpPreparation::ApplyTensor(formatCastOfSelf);
50	- formatCastOfSelf.sizes(), formatCastOfSelf.options(), CalcuOpUtil::get_tensor_npu_format(formatCastOfSelf));	49	+ Tensor mask = OpPreparation::ApplyTensor(formatCastOfSelf, formatCastOfSeed.options());
51	- Tensor mask = at::empty_with_format(
52	- formatCastOfSelf.sizes(), formatCastOfSeed.options(), CalcuOpUtil::get_tensor_npu_format(formatCastOfSelf));
53	-
54	dropout_v2_out_npu(result, mask, formatCastOfSeed, formatCastOfSelf, formatCastOfSeed, p);	50	dropout_v2_out_npu(result, mask, formatCastOfSeed, formatCastOfSelf, formatCastOfSeed, p);
55	NpuUtils::format_fresh_view(seed, formatCastOfSeed);	51	NpuUtils::format_fresh_view(seed, formatCastOfSeed);
56	return std::tuple<Tensor, Tensor, Tensor>(result, mask, seed);	52	return std::tuple<Tensor, Tensor, Tensor>(result, mask, seed);

Msrc/aten/src/ATen/native/npu/EmbeddingKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	#include "c10/npu/OptionsManager.h"	19	#include "c10/npu/OptionsManager.h"
19		20
20	namespace at {	21	namespace at {

Msrc/aten/src/ATen/native/npu/EqKernelNpu.cpp+2-2

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19		19
20	namespace at {	20	namespace at {
21	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/ExpKernelNpu.cpp+2-8

	@@ -34,9 +34,7 @@ Tensor& exp_out_npu(Tensor& result, const Tensor& self) {
34	OpPreparation::CheckOut(	34	OpPreparation::CheckOut(
35	{self},	35	{self},
36	result,	36	result,
37	- CalcuOpUtil::get_tensor_npu_format(self),	37	+ self);
38	- self.scalar_type(),
39	- self.sizes());
40		38
41	OpPipeWithDefinedOut pipe;	39	OpPipeWithDefinedOut pipe;
42	return pipe.CheckMemory({self}, {result})	40	return pipe.CheckMemory({self}, {result})
	@@ -51,11 +49,7 @@ Tensor& exp_npu_(Tensor& self) {
51	}	49	}
52		50
53	Tensor exp_npu(const Tensor& self) {	51	Tensor exp_npu(const Tensor& self) {
54	- // construct the output tensor of the NPU	52	+ Tensor result = OpPreparation::ApplyTensor(self);
55	- Tensor result = at::empty_with_format(
56	- self.sizes(), self.options(), CalcuOpUtil::get_tensor_npu_format(self));
57	-
58	- // calculate the output result of the NPU
59	exp_out_npu_nocheck(result, self);	53	exp_out_npu_nocheck(result, self);
60	return result;	54	return result;
61	}	55	}

Msrc/aten/src/ATen/native/npu/FillKernelNpu.cpp+1-2

	@@ -44,12 +44,11 @@ Tensor& fill_out_npu(Tensor& result, Tensor& self, const Tensor& other) {
44	Tensor& fills_out_npu(Tensor& result, Tensor& self, Scalar value) {	44	Tensor& fills_out_npu(Tensor& result, Tensor& self, Scalar value) {
45	AT_DISPATCH_ALL_TYPES_AND3(kHalf, kBool, kBFloat16, self.scalar_type(), "fills_out_npu", [&]() {	45	AT_DISPATCH_ALL_TYPES_AND3(kHalf, kBool, kBFloat16, self.scalar_type(), "fills_out_npu", [&]() {
46	auto value_converted = value.to<scalar_t>();});	46	auto value_converted = value.to<scalar_t>();});
47	- float scalar = CalcuOpUtil::get_scalar_float_value(value);
48	OpCommand cmd;	47	OpCommand cmd;
49	cmd.Name("Fills")	48	cmd.Name("Fills")
50	.Input(self)	49	.Input(self)
51	.Output(result)	50	.Output(result)
52	- .Attr("value", scalar)	51	+ .Attr("value", value)
53	.Run();	52	.Run();
54		53
55	return result;	54	return result;

Msrc/aten/src/ATen/native/npu/FlipKernelNpu.cpp+1-7

	@@ -12,7 +12,6 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
16	#include "c10/npu/OptionsManager.h"	15	#include "c10/npu/OptionsManager.h"
17	#include "ATen/native/npu/utils/OpAdapter.h"	16	#include "ATen/native/npu/utils/OpAdapter.h"
18		17
	@@ -21,12 +20,7 @@ namespace native {
21	using namespace at::native::npu;	20	using namespace at::native::npu;
22		21
23	Tensor flip_npu(const Tensor& self, IntArrayRef dims){	22	Tensor flip_npu(const Tensor& self, IntArrayRef dims){
24	- // calculate the output size	23	+ Tensor result = OpPreparation::ApplyTensor(self);
25	- auto outputSize = input_same_output_size(self);
26	-
27	- // construct the output tensor of the NPU
28	- Tensor result = at::empty_with_format(outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
29	-
30	SmallVector<int64_t,N> dimVec = array_to_small_vector(dims);	24	SmallVector<int64_t,N> dimVec = array_to_small_vector(dims);
31	if (!c10::npu::OptionsManager::CheckDynamicEnable()) {	25	if (!c10::npu::OptionsManager::CheckDynamicEnable()) {
32	OpCommand cmd;	26	OpCommand cmd;

Msrc/aten/src/ATen/native/npu/FloorDivideKernelNpu.cpp+2-2

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
19		19
20	namespace at {	20	namespace at {
21	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/FloorKernelNpu.cpp+2-8

	@@ -34,9 +34,7 @@ Tensor& floor_out_npu(Tensor& result, const Tensor& self) {
34	OpPreparation::CheckOut(	34	OpPreparation::CheckOut(
35	{self},	35	{self},
36	result,	36	result,
37	- CalcuOpUtil::get_tensor_npu_format(self),	37	+ self);
38	- self.scalar_type(),
39	- self.sizes());
40		38
41	OpPipeWithDefinedOut pipe;	39	OpPipeWithDefinedOut pipe;
42	return pipe.CheckMemory({self}, {result})	40	return pipe.CheckMemory({self}, {result})
	@@ -51,11 +49,7 @@ Tensor& floor_npu_(Tensor& self) {
51	}	49	}
52		50
53	Tensor floor_npu(const Tensor& self) {	51	Tensor floor_npu(const Tensor& self) {
54	- // construct the output tensor of the NPU	52	+ Tensor result = OpPreparation::ApplyTensor(self);
55	- Tensor result = at::empty_with_format(
56	- self.sizes(), self.options(), CalcuOpUtil::get_tensor_npu_format(self));
57	-
58	- // calculate the output result of the NPU
59	floor_out_npu_nocheck(result, self);	53	floor_out_npu_nocheck(result, self);
60	return result;	54	return result;
61	}	55	}

Msrc/aten/src/ATen/native/npu/FmodKernelNpu.cpp+7-27

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -49,9 +48,8 @@ Tensor& fmod_out_npu(Tensor& result, const Tensor& self, const Tensor& other) {
49	auto outputSize = broadcast_ops_npu_output_size(self, other);	48	auto outputSize = broadcast_ops_npu_output_size(self, other);
50	OpPreparation::CheckOut(	49	OpPreparation::CheckOut(
51	{self, other},	50	{self, other},
52	- result,	51	+ result,
53	- CalcuOpUtil::get_tensor_npu_format(self),	52	+ self,
54	- self.scalar_type(),
55	outputSize);	53	outputSize);
56		54
57	fmod_out_npu_nocheck(result, self, other);	55	fmod_out_npu_nocheck(result, self, other);
	@@ -65,10 +63,7 @@ Tensor& fmod_out_npu(Tensor& result, const Tensor& self, Scalar other) {
65	}	63	}
66		64
67	Tensor& fmod_npu_(Tensor& self, Scalar other) {	65	Tensor& fmod_npu_(Tensor& self, Scalar other) {
68	- SmallVector<Tensor, N> inputs = {self};	66	+ OpPreparation::CheckMemory({self}, {self});
69	- SmallVector<Tensor, N> outputs = {self};
70	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
71	-
72	if (!NpuUtils::check_match(&self)) {	67	if (!NpuUtils::check_match(&self)) {
73	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	68	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
74	Tensor result = fmod_out_npu_nocheck(contiguousSelf, contiguousSelf, other);	69	Tensor result = fmod_out_npu_nocheck(contiguousSelf, contiguousSelf, other);
	@@ -81,10 +76,7 @@ Tensor& fmod_npu_(Tensor& self, Scalar other) {
81	}	76	}
82		77
83	Tensor& fmod_npu_(Tensor& self, const Tensor& other) {	78	Tensor& fmod_npu_(Tensor& self, const Tensor& other) {
84	- SmallVector<Tensor, N> inputs = {self, other};	79	+ OpPreparation::CheckMemory({self, other}, {self});
85	- SmallVector<Tensor, N> outputs = {self};
86	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
87	-
88	if (!NpuUtils::check_match(&self)) {	80	if (!NpuUtils::check_match(&self)) {
89	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	81	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
90	Tensor result = fmod_out_npu_nocheck(contiguousSelf, contiguousSelf, other);	82	Tensor result = fmod_out_npu_nocheck(contiguousSelf, contiguousSelf, other);
	@@ -97,26 +89,14 @@ Tensor& fmod_npu_(Tensor& self, const Tensor& other) {
97	}	89	}
98		90
99	Tensor fmod_npu(const Tensor& self, Scalar other) {	91	Tensor fmod_npu(const Tensor& self, Scalar other) {
100	- // calculate the output size	92	+ Tensor result = OpPreparation::ApplyTensor(self);
101	- auto outputSize = input_same_output_size(self);
102	- // construct the output tensor of the NPU
103	- Tensor result = at::empty_with_format(
104	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
105	-
106	- // calculate the output result of the NPU
107	fmod_out_npu_nocheck(result, self, other);	93	fmod_out_npu_nocheck(result, self, other);
108	return result;	94	return result;
109	}	95	}
110		96
111	Tensor fmod_npu(const Tensor& self, const Tensor& other) {	97	Tensor fmod_npu(const Tensor& self, const Tensor& other) {
112	- // calculate the output size
113	auto outputSize = broadcast_ops_npu_output_size(self, other);	98	auto outputSize = broadcast_ops_npu_output_size(self, other);
114	-	99	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
115	- // construct the output tensor of the NPU
116	- Tensor result = at::empty_with_format(
117	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
118	-
119	- // calculate the output result of the NPU
120	fmod_out_npu_nocheck(result, self, other);	100	fmod_out_npu_nocheck(result, self, other);
121	return result;	101	return result;
122	}	102	}

Msrc/aten/src/ATen/native/npu/GatherKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/GeKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/GeluKernelNpu.cpp+4-6

	@@ -11,18 +11,16 @@
11	// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.	11	// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	14	+
15	-#include "ATen/native/npu/utils/OpTemplate.h"	15	+
		16	+#include "ATen/native/npu/utils/OpAdapter.h"
16		17
17	namespace at {	18	namespace at {
18	namespace native {	19	namespace native {
19	using namespace at::native::npu;	20	using namespace at::native::npu;
20		21
21	Tensor gelu_npu(const Tensor& self) {	22	Tensor gelu_npu(const Tensor& self) {
22	- // calculate the output size	23	+ Tensor result = OpPreparation::ApplyTensor(self);
23	- auto outputSize = input_same_output_size(self);
24	- // construct the output tensor of the NPU
25	- Tensor result = at::empty_with_format(outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
26	// calculate the output result of the NPU	24	// calculate the output result of the NPU
27	OpCommand cmd;	25	OpCommand cmd;
28	cmd.Name("Gelu")	26	cmd.Name("Gelu")

Msrc/aten/src/ATen/native/npu/GerKernelNpu.cpp+1-2

	@@ -54,8 +54,7 @@ Tensor& ger_out_npu(Tensor& result, const Tensor& self , const Tensor& vec2) {
54	OpPreparation::CheckOut(	54	OpPreparation::CheckOut(
55	{self},	55	{self},
56	result,	56	result,
57	- CalcuOpUtil::get_tensor_npu_format(self),	57	+ self,
58	- self.scalar_type(),
59	outputSize);	58	outputSize);
60		59
61	OpPipeWithDefinedOut pipe;	60	OpPipeWithDefinedOut pipe;

Msrc/aten/src/ATen/native/npu/GridSampler3dBackwardKernelNpu.cpp+0-1

	@@ -15,7 +15,6 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include <c10/npu/OptionsManager.h>	17	#include <c10/npu/OptionsManager.h>
18	-#include "ATen/native/npu/utils/CalcuOpUtil.h"
19	#include "ATen/native/npu/utils/OpAdapter.h"	18	#include "ATen/native/npu/utils/OpAdapter.h"
20		19
21	namespace at {	20	namespace at {

Msrc/aten/src/ATen/native/npu/GridSampler3dKernelNpu.cpp+0-1

	@@ -15,7 +15,6 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include <c10/npu/OptionsManager.h>	17	#include <c10/npu/OptionsManager.h>
18	-#include "ATen/native/npu/utils/CalcuOpUtil.h"
19	#include "ATen/native/npu/utils/OpAdapter.h"	18	#include "ATen/native/npu/utils/OpAdapter.h"
20		19
21	namespace at {	20	namespace at {

Msrc/aten/src/ATen/native/npu/GtKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/HardsigmoidKernelNpu.cpp+1-4

	@@ -39,10 +39,7 @@ Tensor hardsigmoid_npu(const Tensor& self) {
39	}	39	}
40		40
41	Tensor& hardsigmoid_npu_(Tensor& self) {	41	Tensor& hardsigmoid_npu_(Tensor& self) {
42	- SmallVector<Tensor, N> inputs = {self};	42	+ OpPreparation::CheckMemory({self}, {self});
43	- SmallVector<Tensor, N> outputs = {self};
44	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
45	-
46	if (!NpuUtils::check_match(&self)) {	43	if (!NpuUtils::check_match(&self)) {
47	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	44	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
48	Tensor result = hardsigmoid_out_npu(contiguousSelf, contiguousSelf);	45	Tensor result = hardsigmoid_out_npu(contiguousSelf, contiguousSelf);

Msrc/aten/src/ATen/native/npu/HardtanhBackwardKernelNpu.cpp+3-11

	@@ -27,15 +27,13 @@ Tensor& hardtanh_backward_out_npu(
27	const Tensor& self,	27	const Tensor& self,
28	Scalar min_val,	28	Scalar min_val,
29	Scalar max_val) {	29	Scalar max_val) {
30	- float max_value = CalcuOpUtil::get_scalar_float_value(max_val);
31	- float min_value = CalcuOpUtil::get_scalar_float_value(min_val);
32	OpCommand cmd;	30	OpCommand cmd;
33	cmd.Name("HardtanhGrad")	31	cmd.Name("HardtanhGrad")
34	.Input(self)	32	.Input(self)
35	.Input(grad_output)	33	.Input(grad_output)
36	.Output(grad_input)	34	.Output(grad_input)
37	- .Attr("max_val", max_value)	35	+ .Attr("max_val", max_val)
38	- .Attr("min_val", min_value)	36	+ .Attr("min_val", min_val)
39	.Run();	37	.Run();
40		38
41	return grad_input;	39	return grad_input;
	@@ -46,13 +44,7 @@ Tensor hardtanh_backward_npu(
46	const Tensor& self,	44	const Tensor& self,
47	Scalar min_val,	45	Scalar min_val,
48	Scalar max_val) {	46	Scalar max_val) {
49	- // calculate the output size	47	+ Tensor grad_input = OpPreparation::ApplyTensor(self);
50	- auto outputSize = input_same_output_size(self);
51	-
52	- // construct the output tensor of the NPU
53	- Tensor grad_input = at::empty_with_format(
54	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
55	-
56	// calculate the output result of the NPU	48	// calculate the output result of the NPU
57	hardtanh_backward_out_npu(grad_input, grad_output, self, min_val, max_val);	49	hardtanh_backward_out_npu(grad_input, grad_output, self, min_val, max_val);
58		50

Msrc/aten/src/ATen/native/npu/HardtanhKernelNpu.cpp+3-15

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -38,24 +37,13 @@ Tensor& hardtanh_out_npu(
38	}	37	}
39		38
40	Tensor hardtanh_npu(const Tensor& self, Scalar min, Scalar max) {	39	Tensor hardtanh_npu(const Tensor& self, Scalar min, Scalar max) {
41	- // calculate the output size	40	+ Tensor result = OpPreparation::ApplyTensor(self);
42	- auto outputSize = input_same_output_size(self);
43	-
44	- // construct the output tensor of the NPU
45	- Tensor result = at::empty_with_format(
46	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
47	-
48	- // calculate the output result of the NPU
49	hardtanh_out_npu(result, self, min, max);	41	hardtanh_out_npu(result, self, min, max);
50	-
51	return result;	42	return result;
52	}	43	}
53		44
54	Tensor& hardtanh_npu_(Tensor& self, Scalar min, Scalar max) {	45	Tensor& hardtanh_npu_(Tensor& self, Scalar min, Scalar max) {
55	- SmallVector<Tensor, N> inputs = {self};	46	+ OpPreparation::CheckMemory({self}, {self});
56	- SmallVector<Tensor, N> outputs = {self};
57	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
58	-
59	if (!NpuUtils::check_match(&self)) {	47	if (!NpuUtils::check_match(&self)) {
60	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	48	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
61	Tensor result = hardtanh_out_npu(contiguousSelf, contiguousSelf, min, max);	49	Tensor result = hardtanh_out_npu(contiguousSelf, contiguousSelf, min, max);

Msrc/aten/src/ATen/native/npu/IfmrKernelNpu.cpp+1-3

	@@ -14,9 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/CalcuOpUtil.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
19	-#include "ATen/native/npu/utils/OpTemplate.h"
20		18
21	namespace at {	19	namespace at {
22	namespace native {	20	namespace native {

Msrc/aten/src/ATen/native/npu/Im2colBackwardKernelNpu.cpp+1-2

	@@ -71,8 +71,7 @@ Tensor& im2col_backward_out_npu(
71	OpPreparation::CheckOut(	71	OpPreparation::CheckOut(
72	{grad_output},	72	{grad_output},
73	grad_input,	73	grad_input,
74	- CalcuOpUtil::get_tensor_npu_format(grad_output),	74	+ grad_output,
75	- grad_output.scalar_type(),
76	outputSize);	75	outputSize);
77		76
78	OpPipeWithDefinedOut pipe;	77	OpPipeWithDefinedOut pipe;

Msrc/aten/src/ATen/native/npu/Im2colKernelNpu.cpp+2-6

	@@ -124,8 +124,7 @@ Tensor& im2col_out_npu(Tensor& result, const Tensor &self, IntArrayRef kernel_si
124	OpPreparation::CheckOut(	124	OpPreparation::CheckOut(
125	{self},	125	{self},
126	result,	126	result,
127	- CalcuOpUtil::get_tensor_npu_format(self),	127	+ self,
128	- self.scalar_type(),
129	image_to_col_npu_output_size(self, kernel_size, stride, dilation, padding));	128	image_to_col_npu_output_size(self, kernel_size, stride, dilation, padding));
130		129
131	OpPipeWithDefinedOut pipe;	130	OpPipeWithDefinedOut pipe;
	@@ -139,10 +138,7 @@ Tensor im2col_npu(const Tensor &self, IntArrayRef kernel_size, IntArrayRef dilat
139	// calculate the output size	138	// calculate the output size
140	auto outputSize =	139	auto outputSize =
141	image_to_col_npu_output_size(self, kernel_size, stride, dilation, padding);	140	image_to_col_npu_output_size(self, kernel_size, stride, dilation, padding);
142	-	141	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
143	- Tensor result = at::empty_with_format(
144	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
145	-
146	im2col_out_npu(result, self, kernel_size, dilation, padding, stride);	142	im2col_out_npu(result, self, kernel_size, dilation, padding, stride);
147		143
148	return result;	144	return result;

Msrc/aten/src/ATen/native/npu/IndexAddKernelNpu.cpp+2-4

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include<ATen/NamedTensorUtils.h>
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {
	@@ -51,10 +52,7 @@ Tensor& index_add_npu_(
51	int64_t dim,	52	int64_t dim,
52	const Tensor& index,	53	const Tensor& index,
53	const Tensor& source) {	54	const Tensor& source) {
54	- SmallVector<Tensor, N> inputs = {self, index, source};	55	+ OpPreparation::CheckMemory({self, index, source}, {self});
55	- SmallVector<Tensor, N> outputs = {self};
56	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
57	-
58	if (!NpuUtils::check_match(&self)) {	56	if (!NpuUtils::check_match(&self)) {
59	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	57	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
60	Tensor result = index_add_out_npu(contiguousSelf, contiguousSelf, dim, index, source);	58	Tensor result = index_add_out_npu(contiguousSelf, contiguousSelf, dim, index, source);

Msrc/aten/src/ATen/native/npu/IndexFillDKernelNpu.cpp+1-0

	@@ -13,6 +13,7 @@
13	// limitations under the License.	13	// limitations under the License.
14		14
15	#include "ATen/native/npu/utils/OpAdapter.h"	15	#include "ATen/native/npu/utils/OpAdapter.h"
		16	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
16	#include <vector>	17	#include <vector>
17		18
18	namespace at{	19	namespace at{

Msrc/aten/src/ATen/native/npu/IndexPutKernelNpu.cpp+3-6

	@@ -14,6 +14,7 @@
14	// limitations under the License.	14	// limitations under the License.
15		15
16	#include "ATen/native/npu/utils/OpAdapter.h"	16	#include "ATen/native/npu/utils/OpAdapter.h"
		17	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
17		18
18	namespace at {	19	namespace at {
19	namespace native {	20	namespace native {
	@@ -37,10 +38,9 @@ Tensor& index_put_nocheck(
37	}	38	}
38	}	39	}
39		40
40	- Tensor masksTensor = CalcuOpUtil::copy_tensor_host_to_device(	41	+ auto masksTensor = CalcuOpUtil::copy_tensor_host_to_device(
41	from_blob(masks.data(), {masks.size()}, dtype(ScalarType::Long)));	42	from_blob(masks.data(), {masks.size()}, dtype(ScalarType::Long)));
42		43
43	-
44	OpCommand cmd;	44	OpCommand cmd;
45	cmd.Name("IndexPut")	45	cmd.Name("IndexPut")
46	.Input(self)	46	.Input(self)
	@@ -78,10 +78,7 @@ Tensor& _index_put_impl_npu_(
78	const Tensor& value,	78	const Tensor& value,
79	const bool accumulate,	79	const bool accumulate,
80	const bool unsafe) {	80	const bool unsafe) {
81	- SmallVector<Tensor, N> inputs = {self};	81	+ OpPreparation::CheckMemory({self}, {self});
82	- SmallVector<Tensor, N> outputs = {self};
83	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
84	-
85	OpPreparation::CastBackToOriFormat(self);	82	OpPreparation::CastBackToOriFormat(self);
86		83
87	Tensor valueCopy = value;	84	Tensor valueCopy = value;

Msrc/aten/src/ATen/native/npu/IndexSelectKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	#include "c10/npu/OptionsManager.h"	19	#include "c10/npu/OptionsManager.h"
19		20
20	namespace at {	21	namespace at {

Msrc/aten/src/ATen/native/npu/IndexingKernelNpu.cpp+2-5

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -56,9 +55,7 @@ Tensor indexing_npu(
56	outputSize.emplace_back((end[i] + strides[i] - 1 - begin[i]) / strides[i]);	55	outputSize.emplace_back((end[i] + strides[i] - 1 - begin[i]) / strides[i]);
57	}	56	}
58	// construct the output tensor of the NPU	57	// construct the output tensor of the NPU
59	- Tensor result = at::empty_with_format(	58	+ Tensor result = OpPreparation::ApplyTensor(self, outputSize);
60	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
61	-
62	indexing_out_npu(result, self, begin, end, strides);	59	indexing_out_npu(result, self, begin, end, strides);
63		60
64	return result;	61	return result;

Dsrc/aten/src/ATen/native/npu/InstanceNormKernelNpu.cpp+0-95

	@@ -1,95 +0,0 @@
1	-// Copyright (c) 2020, Huawei Technologies.All rights reserved.
2	-//
3	-// Licensed under the BSD 3-Clause License (the "License");
4	-// you may not use this file except in compliance with the License.
5	-// You may obtain a copy of the License at
6	-//
7	-// https://opensource.org/licenses/BSD-3-Clause
8	-//
9	-// Unless required by applicable law or agreed to in writing, software
10	-// distributed under the License is distributed on an "AS IS" BASIS,
11	-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12	-// See the License for the specific language governing permissions and
13	-// limitations under the License.
14	-
15	-#include <c10/npu/NPUCachingAllocator.h>
16	-#include "ATen/native/npu/utils/CalcuOpUtil.h"
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"
18	-#include "ATen/native/npu/utils/NpuUtils.h"
19	-
20	-namespace at {
21	-namespace native {
22	-using namespace at::native::npu;
23	-
24	-SmallVector<NPUTensorDesc, N> instance_norm_npu_input(const SmallVector<Tensor, N>& inputTensor) {
25	- return CalcuOpUtil::create_npu_input_tensor_desc(inputTensor);
26	-}
27	-
28	-SmallVector<NPUTensorDesc, N> instance_norm_npu_output(const Tensor& result) {
29	- return CalcuOpUtil::create_npu_output_tensor_desc({result});
30	-}
31	-
32	-SmallVector<NPUAttrDesc, N> instance_norm_npu_attr(bool use_input_stats, double momentum, double eps) {
33	- NPUAttrDesc npuAttrStats = NPUAttrDesc("use_input_stats", use_input_stats);
34	- NPUAttrDesc npuAttrMomentum = NPUAttrDesc("momentum", static_cast<float>(momentum));
35	- NPUAttrDesc npuAttrEpsilon = NPUAttrDesc("eps", static_cast<float>(eps));
36	- SmallVector<NPUAttrDesc, N> attrs = {npuAttrStats, npuAttrMomentum, npuAttrEpsilon};
37	- return attrs;
38	-}
39	-
40	-Tensor& instance_norm_out_npu(
41	- Tensor& result,
42	- const Tensor& self,
43	- const Tensor& weight,
44	- const Tensor& bias,
45	- const Tensor& running_mean,
46	- const Tensor& running_var,
47	- bool use_input_stats,
48	- double momentum,
49	- double eps) {
50	- // constructs the input and output NPUTensorDesc
51	- auto inputs = instance_norm_npu_input(
52	- {self, weight, bias, running_mean, running_var});
53	- auto outputs = instance_norm_npu_output(result);
54	-
55	- // constructs the attr of the NPUAttrDesc
56	- auto attrs = instance_norm_npu_attr(use_input_stats, momentum, eps);
57	-
58	- // executing the NPU operator
59	- CalcuOpUtil::execute_npu_operate("InstanceNorm", inputs, outputs, attrs);
60	-
61	- return result;
62	-}
63	-
64	-Tensor instance_norm_npu(
65	- const Tensor& self,
66	- const Tensor& weight,
67	- const Tensor& bias,
68	- const Tensor& running_mean,
69	- const Tensor& running_var,
70	- bool use_input_stats,
71	- double momentum,
72	- double eps,
73	- bool cudnn_enabled) {
74	- TORCH_CHECK(use_input_stats \|\| (running_mean.defined() && running_var.defined()),
75	- "Expected running_mean and running_var to be defined when use_input_stats is false");
76	- Tensor result = at::empty_with_format(self.sizes(), self.options(),
77	- CalcuOpUtil::get_tensor_npu_format(self));
78	-
79	- // calculate the output result of the NPU
80	- instance_norm_out_npu(
81	- result,
82	- self,
83	- weight,
84	- bias,
85	- running_mean,
86	- running_var,
87	- use_input_stats,
88	- momentum,
89	- eps);
90	-
91	- return result;
92	-}
93	-
94	-} // namespace native
95	-} // namespace at

Msrc/aten/src/ATen/native/npu/InverseKernelNpu.cpp+2-4

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -36,8 +35,7 @@ Tensor& inverse_out_npu(
36	}	35	}
37		36
38	Tensor inverse_npu(const Tensor& self) {	37	Tensor inverse_npu(const Tensor& self) {
39	- Tensor result = at::empty_with_format(	38	+ Tensor result = OpPreparation::ApplyTensor(self);
40	- self.sizes(), self.options(), CalcuOpUtil::get_tensor_npu_format(self));
41		39
42	inverse_out_npu(result, self);	40	inverse_out_npu(result, self);
43		41

Msrc/aten/src/ATen/native/npu/L1LossBackwardKernelNpu.cpp+1-0

	@@ -13,6 +13,7 @@
13	// limitations under the License.	13	// limitations under the License.
14		14
15	#include "ATen/native/npu/utils/OpAdapter.h"	15	#include "ATen/native/npu/utils/OpAdapter.h"
		16	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
16		17
17	namespace at {	18	namespace at {
18	namespace native {	19	namespace native {

Msrc/aten/src/ATen/native/npu/L1lossKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/LeKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/LeakyReluBackwardKernelNpu.cpp+3-12

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -27,13 +26,12 @@ Tensor leaky_relu_backward_out_npu(
27	const Tensor& self,	26	const Tensor& self,
28	Scalar negval,	27	Scalar negval,
29	bool is_result) {	28	bool is_result) {
30	- float negvalValue = CalcuOpUtil::get_scalar_float_value(negval);
31	OpCommand cmd;	29	OpCommand cmd;
32	cmd.Name("LeakyReluGrad")	30	cmd.Name("LeakyReluGrad")
33	.Input(grad_output)	31	.Input(grad_output)
34	.Input(self)	32	.Input(self)
35	.Output(result)	33	.Output(result)
36	- .Attr("negative_slope", negvalValue)	34	+ .Attr("negative_slope", negval)
37	.Run();	35	.Run();
38	return result;	36	return result;
39	}	37	}
	@@ -43,14 +41,7 @@ Tensor leaky_relu_backward_npu(
43	const Tensor& self,	41	const Tensor& self,
44	Scalar negval,	42	Scalar negval,
45	bool is_result) {	43	bool is_result) {
46	- // calculate the output size	44	+ Tensor result = OpPreparation::ApplyTensor(self);
47	- auto outputSize = input_same_output_size(self);
48	-
49	- // construct the output tensor of the NPU
50	- Tensor result = at::empty_with_format(
51	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
52	-
53	- // calculate the output result of the NPU
54	leaky_relu_backward_out_npu(result, grad_output, self, negval, is_result);	45	leaky_relu_backward_out_npu(result, grad_output, self, negval, is_result);
55	return result;	46	return result;
56	}	47	}

Msrc/aten/src/ATen/native/npu/LeakyReluKernelNpu.cpp+3-13

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -33,23 +32,14 @@ Tensor& leaky_relu_out_npu(Tensor& result, const Tensor& self, Scalar negval) {
33	}	32	}
34		33
35	Tensor leaky_relu_npu(const Tensor& self, Scalar negval) {	34	Tensor leaky_relu_npu(const Tensor& self, Scalar negval) {
36	- // calculate the output size	35	+ Tensor result = OpPreparation::ApplyTensor(self);
37	- auto outputSize = input_same_output_size(self);
38	-
39	- // construct the output tensor of the NPU
40	- Tensor result = at::empty_with_format(
41	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
42	-
43	// calculate the output result of the NPU	36	// calculate the output result of the NPU
44	leaky_relu_out_npu(result, self, negval);	37	leaky_relu_out_npu(result, self, negval);
45	return result;	38	return result;
46	}	39	}
47		40
48	Tensor& leaky_relu_npu_(Tensor& self, Scalar neg_val) {	41	Tensor& leaky_relu_npu_(Tensor& self, Scalar neg_val) {
49	- SmallVector<Tensor, N> inputs = {self};	42	+ OpPreparation::CheckMemory({self}, {self});
50	- SmallVector<Tensor, N> outputs = {self};
51	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
52	-
53	if (!NpuUtils::check_match(&self)) {	43	if (!NpuUtils::check_match(&self)) {
54	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	44	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
55	Tensor result = leaky_relu_out_npu(contiguousSelf, contiguousSelf, neg_val);	45	Tensor result = leaky_relu_out_npu(contiguousSelf, contiguousSelf, neg_val);

Msrc/aten/src/ATen/native/npu/LinspaceKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Asrc/aten/src/ATen/native/npu/Log10KernelNpu.cpp+65-0

@@ -0,0 +1,65 @@
	1	+// Copyright (c) 2020 Huawei Technologies Co., Ltd
	2	+// Copyright (c) 2019, Facebook CORPORATION.
	3	+// All rights reserved.
	4	+//
	5	+// Licensed under the BSD 3-Clause License (the "License");
	6	+// you may not use this file except in compliance with the License.
	7	+// You may obtain a copy of the License at
	8	+//
	9	+// https://opensource.org/licenses/BSD-3-Clause
	10	+//
	11	+// Unless required by applicable law or agreed to in writing, software
	12	+// distributed under the License is distributed on an "AS IS" BASIS,
	13	+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	14	+// See the License for the specific language governing permissions and
	15	+// limitations under the License.
	16	+
	17	+#include "ATen/native/npu/utils/OpAdapter.h"
	18	+
	19	+namespace at {
	20	+namespace native {
	21	+using namespace at::native::npu;
	22	+
	23	+Tensor& log10_out_npu_nocheck(Tensor& result, const Tensor& self) {
	24	+ OpCommand cmd;
	25	+ cmd.Name("Log")
	26	+ .Input(self)
	27	+ .Output(result)
	28	+ .Attr("base", (float)10.0)
	29	+ .Attr("scale", (float)1.0)
	30	+ .Attr("shift", (float)0.0)
	31	+ .Run();
	32	+
	33	+ return result;
	34	+}
	35	+
	36	+Tensor& log10_out_npu(Tensor& result, const Tensor& self) {
	37	+ OpPreparation::CheckOut(
	38	+ {self},
	39	+ result,
	40	+ self);
	41	+
	42	+ OpPipeWithDefinedOut pipe;
	43	+ return pipe.CheckMemory({self}, {result})
	44	+ .Func([&self](Tensor& result){log10_out_npu_nocheck(result, self);})
	45	+ .Call(result);
	46	+}
	47	+
	48	+Tensor log10_npu(const Tensor& self) {
	49	+ // construct the output tensor of the NPU
	50	+ Tensor result = OpPreparation::ApplyTensor(self);
	51	+
	52	+ // calculate the output result of the NPU
	53	+ log10_out_npu_nocheck(result, self);
	54	+
	55	+ return result;
	56	+}
	57	+
	58	+Tensor& log10_npu_(Tensor& self) {
	59	+ log10_out_npu(self, self);
	60	+
	61	+ return self;
	62	+}
	63	+
	64	+} // namespace native
	65	+} // namespace at

Msrc/aten/src/ATen/native/npu/Log1pKernelNpu.cpp+2-10

	@@ -29,13 +29,7 @@ Tensor& log1p_out_npu(Tensor& result, const Tensor& self){
29	}	29	}
30		30
31	Tensor log1p_npu(const Tensor& self) {	31	Tensor log1p_npu(const Tensor& self) {
32	- //calculate the output size	32	+ Tensor result = OpPreparation::ApplyTensor(self);
33	- auto outputSize = input_same_output_size(self);
34	-
35	- //construct the output tensor of the NPU
36	- Tensor result = at::empty_with_format(outputSize,
37	- self.options(), CalcuOpUtil::get_tensor_npu_format(self));
38	-
39	//calculate the output result of the NPU	33	//calculate the output result of the NPU
40	log1p_out_npu(result, self);	34	log1p_out_npu(result, self);
41		35
	@@ -43,9 +37,7 @@ Tensor log1p_npu(const Tensor& self) {
43	}	37	}
44		38
45	Tensor& log1p_npu_(Tensor& self) {	39	Tensor& log1p_npu_(Tensor& self) {
46	- SmallVector<Tensor, N> inputs = {self};	40	+ OpPreparation::CheckMemory({self}, {self});
47	- SmallVector<Tensor, N> outputs = {self};
48	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
49		41
50	if (!NpuUtils::check_match(&self)) {	42	if (!NpuUtils::check_match(&self)) {
51	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	43	Tensor contiguousSelf = NpuUtils::format_contiguous(self);

Msrc/aten/src/ATen/native/npu/Log2KernelNpu.cpp+3-15

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -35,24 +34,13 @@ Tensor& log2_out_npu(Tensor& result, const Tensor& self) {
35	}	34	}
36		35
37	Tensor log2_npu(const Tensor& self) {	36	Tensor log2_npu(const Tensor& self) {
38	- // calculate the output size	37	+ Tensor result = OpPreparation::ApplyTensor(self);
39	- auto outputSize = input_same_output_size(self);
40	-
41	- // construct the output tensor of the NPU
42	- Tensor result = at::empty_with_format(
43	- outputSize, self.options(), CalcuOpUtil::get_tensor_npu_format(self));
44	-
45	- // calculate the output result of the NPU
46	log2_out_npu(result, self);	38	log2_out_npu(result, self);
47	-
48	return result;	39	return result;
49	}	40	}
50		41
51	Tensor& log2_npu_(Tensor& self) {	42	Tensor& log2_npu_(Tensor& self) {
52	- SmallVector<Tensor, N> inputs = {self};	43	+ OpPreparation::CheckMemory({self}, {self});
53	- SmallVector<Tensor, N> outputs = {self};
54	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
55	-
56	if (!NpuUtils::check_match(&self)) {	44	if (!NpuUtils::check_match(&self)) {
57	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	45	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
58	Tensor result = log2_out_npu(contiguousSelf, contiguousSelf);	46	Tensor result = log2_out_npu(contiguousSelf, contiguousSelf);

Msrc/aten/src/ATen/native/npu/LogKernelNpu.cpp+4-9

	@@ -21,16 +21,13 @@ namespace native {
21	using namespace at::native::npu;	21	using namespace at::native::npu;
22		22
23	Tensor& log_out_npu_nocheck(Tensor& result, const Tensor& self) {	23	Tensor& log_out_npu_nocheck(Tensor& result, const Tensor& self) {
24	- float baseValue = CalcuOpUtil::get_scalar_float_value(-1);
25	- float scaleValue = CalcuOpUtil::get_scalar_float_value(1);
26	- float shiftValue = CalcuOpUtil::get_scalar_float_value(0);
27	OpCommand cmd;	24	OpCommand cmd;
28	cmd.Name("Log")	25	cmd.Name("Log")
29	.Input(self)	26	.Input(self)
30	.Output(result)	27	.Output(result)
31	- .Attr("base", baseValue)	28	+ .Attr("base", (float)-1)
32	- .Attr("scale", scaleValue)	29	+ .Attr("scale", (float)1)
33	- .Attr("shift", shiftValue)	30	+ .Attr("shift", (float)0)
34	.Run();	31	.Run();
35		32
36	return result;	33	return result;
	@@ -40,9 +37,7 @@ Tensor& log_out_npu(Tensor& result, const Tensor& self) {
40	OpPreparation::CheckOut(	37	OpPreparation::CheckOut(
41	{self},	38	{self},
42	result,	39	result,
43	- CalcuOpUtil::get_tensor_npu_format(self),	40	+ self);
44	- self.scalar_type(),
45	- self.sizes());
46		41
47	OpPipeWithDefinedOut pipe;	42	OpPipeWithDefinedOut pipe;
48	return pipe.CheckMemory({self}, {result})	43	return pipe.CheckMemory({self}, {result})

Msrc/aten/src/ATen/native/npu/LogSigmoidKernelNpu.cpp+2-6

	@@ -12,8 +12,7 @@
12	// See the License for the specific language governing permissions and	12	// See the License for the specific language governing permissions and
13	// limitations under the License.	13	// limitations under the License.
14		14
15	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	15	+#include "ATen/native/npu/utils/OpAdapter.h"
16	-#include "ATen/native/npu/utils/OpTemplate.h"
17		16
18	namespace at {	17	namespace at {
19	namespace native {	18	namespace native {
	@@ -33,11 +32,8 @@ tuple<Tensor&, Tensor&> log_sigmoid_forward_out_npu(
33	}	32	}
34		33
35	tuple<Tensor, Tensor> log_sigmoid_forward_npu(const Tensor& self) {	34	tuple<Tensor, Tensor> log_sigmoid_forward_npu(const Tensor& self) {
36	- // construct the output tensor of the NPU	35	+ Tensor output = OpPreparation::ApplyTensor(self);
37	- Tensor output = at::empty_with_format(
38	- self.sizes(), self.options(), CalcuOpUtil::get_tensor_npu_format(self));
39	Tensor buffer = at::empty({0}, self.options());	36	Tensor buffer = at::empty({0}, self.options());
40	-
41	// calculate the output result of the NPU	37	// calculate the output result of the NPU
42	log_sigmoid_forward_out_npu(output, buffer, self);	38	log_sigmoid_forward_out_npu(output, buffer, self);
43	return tuple<Tensor, Tensor>(output, buffer);	39	return tuple<Tensor, Tensor>(output, buffer);

Msrc/aten/src/ATen/native/npu/LogSoftmaxBackwardKernelNpu.cpp+2-11

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -27,16 +26,8 @@ Tensor _log_softmax_backward_npu(
27	int64_t dim,	26	int64_t dim,
28	const Tensor& self) {	27	const Tensor& self) {
29	SmallVector<int64_t, N> dimList = {dim};	28	SmallVector<int64_t, N> dimList = {dim};
30	- // calculate the output size	29	+ Tensor grad_input = OpPreparation::ApplyTensor(grad_output);
31	- auto outputSize = input_same_output_size(grad_output);
32		30
33	- // construct the output tensor of the NPU
34	- Tensor grad_input = at::empty_with_format(
35	- outputSize,
36	- grad_output.options(),
37	- CalcuOpUtil::get_tensor_npu_format(grad_output));
38	-
39	- // calculate the output result of the NPU
40	OpCommand cmd;	31	OpCommand cmd;
41	cmd.Name("LogSoftmaxGrad")	32	cmd.Name("LogSoftmaxGrad")
42	.Input(grad_output)	33	.Input(grad_output)

Msrc/aten/src/ATen/native/npu/LogSoftmaxKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include<ATen/NamedTensorUtils.h>
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/LogSpaceKernelNpu.cpp+2-4

	@@ -40,8 +40,6 @@ Tensor& logspace_out_npu(
40	inputs = at::arange(0, steps, at::device(at::kNPU).dtype(at::kFloat));	40	inputs = at::arange(0, steps, at::device(at::kNPU).dtype(at::kFloat));
41	}	41	}
42		42
43	- float startAttr = CalcuOpUtil::get_scalar_float_value(start);
44	- float endAttr = CalcuOpUtil::get_scalar_float_value(end);
45	int64_t dtype = 0;	43	int64_t dtype = 0;
46	if (result.scalar_type() == at::ScalarType::Half) {	44	if (result.scalar_type() == at::ScalarType::Half) {
47	dtype = 0;	45	dtype = 0;
	@@ -55,8 +53,8 @@ Tensor& logspace_out_npu(
55	cmd.Name("LogSpaceD")	53	cmd.Name("LogSpaceD")
56	.Input(inputs)	54	.Input(inputs)
57	.Output(result)	55	.Output(result)
58	- .Attr("start", startAttr)	56	+ .Attr("start", start)
59	- .Attr("end", endAttr)	57	+ .Attr("end", end)
60	.Attr("steps", steps)	58	.Attr("steps", steps)
61	.Attr("base", static_cast<float>(base))	59	.Attr("base", static_cast<float>(base))
62	.Attr("dtype", dtype)	60	.Attr("dtype", dtype)

Msrc/aten/src/ATen/native/npu/LogicalNotKernelNpu.cpp+3-11

	@@ -14,8 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"
19		18
20	namespace at {	19	namespace at {
21	namespace native {	20	namespace native {
	@@ -49,15 +48,8 @@ Tensor logical_not_npu(const Tensor& self) {
49	}	48	}
50		49
51	Tensor& logical_not_npu_(Tensor& self) {	50	Tensor& logical_not_npu_(Tensor& self) {
52	- SmallVector<Tensor, N> inputs = {self};	51	+ OpPreparation::CheckMemory({self}, {self});
53	- SmallVector<Tensor, N> outputs = {self};	52	+ Tensor result = OpPreparation::ApplyTensor(self, self.options().dtype(ScalarType::Byte));
54	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
55	-
56	- Tensor result = at::empty_with_format(
57	- self.sizes(),
58	- self.options().dtype(ScalarType::Byte),
59	- CalcuOpUtil::get_tensor_npu_format(self));
60	-
61	if (!NpuUtils::check_match(&self)) {	53	if (!NpuUtils::check_match(&self)) {
62	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	54	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
63	logical_not_out_npu(result, contiguousSelf);	55	logical_not_out_npu(result, contiguousSelf);

Msrc/aten/src/ATen/native/npu/LstmKernelNpu.cpp+2-2

	@@ -14,8 +14,8 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/KernelNpuOutputSize.h"	17	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18	-#include "ATen/native/npu/utils/OpTemplate.h"	18	+#include "ATen/native/npu/utils/OpAdapter.h"
19		19
20	namespace at {	20	namespace at {
21	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/LtKernelNpu.cpp+1-0

	@@ -15,6 +15,7 @@
15	// limitations under the License.	15	// limitations under the License.
16		16
17	#include "ATen/native/npu/utils/OpAdapter.h"	17	#include "ATen/native/npu/utils/OpAdapter.h"
		18	+#include "ATen/native/npu/utils/CalcuOpUtil.h"
18		19
19	namespace at {	20	namespace at {
20	namespace native {	21	namespace native {

Msrc/aten/src/ATen/native/npu/MaskedFillKernelNpu.cpp+3-9

	@@ -14,7 +14,7 @@
14	// See the License for the specific language governing permissions and	14	// See the License for the specific language governing permissions and
15	// limitations under the License.	15	// limitations under the License.
16		16
17	-#include "ATen/native/npu/utils/OpTemplate.h"	17	+#include "ATen/native/npu/utils/OpAdapter.h"
18		18
19	namespace at {	19	namespace at {
20	namespace native {	20	namespace native {
	@@ -80,10 +80,7 @@ Tensor& masked_fill_out_npu(Tensor& result, const Tensor& self, const Tensor& ma
80	}	80	}
81		81
82	Tensor& masked_fill_npu_(Tensor& self, const Tensor& mask, const Tensor& value) {	82	Tensor& masked_fill_npu_(Tensor& self, const Tensor& mask, const Tensor& value) {
83	- SmallVector<Tensor, N> inputs = {self, mask, value};	83	+ OpPreparation::CheckMemory({self, mask, value}, {self});
84	- SmallVector<Tensor, N> outputs = {self};
85	-
86	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
87	if (!NpuUtils::check_match(&self)) {	84	if (!NpuUtils::check_match(&self)) {
88	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	85	Tensor contiguousSelf = NpuUtils::format_contiguous(self);
89	Tensor result = masked_fill_out_npu(contiguousSelf, contiguousSelf, mask, value);	86	Tensor result = masked_fill_out_npu(contiguousSelf, contiguousSelf, mask, value);
	@@ -95,10 +92,7 @@ Tensor& masked_fill_npu_(Tensor& self, const Tensor& mask, const Tensor& value)
95	}	92	}
96		93
97	Tensor& masked_fill_npu_(Tensor& self, const Tensor& mask, Scalar value) {	94	Tensor& masked_fill_npu_(Tensor& self, const Tensor& mask, Scalar value) {
98	- SmallVector<Tensor, N> inputs = {self, mask};	95	+ OpPreparation::CheckMemory({self, mask}, {self});
99	- SmallVector<Tensor, N> outputs = {self};
100	-
101	- CalcuOpUtil::check_memory_over_laps(inputs, outputs);
102		96
103	if (!NpuUtils::check_match(&self)) {	97	if (!NpuUtils::check_match(&self)) {
104	Tensor contiguousSelf = NpuUtils::format_contiguous(self);	98	Tensor contiguousSelf = NpuUtils::format_contiguous(self);

Asrc/aten/src/ATen/native/npu/MaskedFillRangeKernelNpu.cpp+74-0

@@ -0,0 +1,74 @@
	1	+// Copyright (c) 2020 Huawei Technologies Co., Ltd
	2	+// Copyright (c) 2019, Facebook CORPORATION.
	3	+// All rights reserved.
	4	+//
	5	+// Licensed under the BSD 3-Clause License (the "License");
	6	+// you may not use this file except in compliance with the License.
	7	+// You may obtain a copy of the License at
	8	+//
	9	+// https://opensource.org/licenses/BSD-3-Clause
	10	+//
	11	+// Unless required by applicable law or agreed to in writing, software
	12	+// distributed under the License is distributed on an "AS IS" BASIS,
	13	+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	14	+// See the License for the specific language governing permissions and
	15	+// limitations under the License.
	16	+
	17	+#include "ATen/native/npu/utils/OpAdapter.h"
	18	+
	19	+namespace at {
	20	+namespace native {
	21	+using namespace at::native::npu;
	22	+
	23	+void mask_fill_range_check(
	24	+ const Tensor& self,
	25	+ const Tensor& start,
	26	+ const Tensor& end,
	27	+ const Tensor& value,
	28	+ int64_t axis){
	29	+ int64_t x_dim = self.dim();
	30	+ int64_t min = -x_dim;
	31	+ int64_t max = x_dim - 1;
	32	+ TORCH_CHECK(
	33	+ !(axis < min \|\| axis > max),
	34	+ "axis overfloaw the range, expected in range [",
	35	+ -x_dim,
	36	+ " ",
	37	+ x_dim - 1,
	38	+ "] ");
	39	+ TORCH_CHECK(
	40	+ start.ndimension() == 2 && start.sizes() == end.sizes(),
	41	+ "Expected noempty 2D start tensor and start' sizes() should be equal end's sizes() ");
	42	+ TORCH_CHECK(
	43	+ start.size(0) == value.size(0),
	44	+ "Expected value.length equal start loop num ");
	45	+ TORCH_CHECK(
	46	+ self.scalar_type() == value.scalar_type(),
	47	+ "value dtype should be equal self dtype !, but value dtype is ",
	48	+ value.scalar_type(),
	49	+ " and self dtype is ",
	50	+ self.scalar_type());
	51	+}
	52	+
	53	+Tensor masked_fill_range_npu(
	54	+ const Tensor& self,
	55	+ const Tensor& start,
	56	+ const Tensor& end,
	57	+ const Tensor& value,
	58	+ int64_t axis){
	59	+ mask_fill_range_check(self, start, end, value, axis);
	60	+ Tensor result = OpPreparation::ApplyTensor(self);
	61	+ OpCommand cmd;
	62	+ cmd.Name("MaskedFillRange")
	63	+ .Input(self)
	64	+ .Input(start)
	65	+ .Input(end)
	66	+ .Input(value)
	67	+ .Output(result)
	68	+ .Attr("axis", axis)
	69	+ .Run();
	70	+ return result;
	71	+}
	72	+
	73	+}
	74	+}

Msrc/aten/src/ATen/native/npu/MaxKernelNpu.cpp+1-1

Msrc/aten/src/ATen/native/npu/MaxV1KernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/MeanKernelNpu.cpp+1-1

Msrc/aten/src/ATen/native/npu/MedianKernelNpu.cpp+1-1

Msrc/aten/src/ATen/native/npu/MinKernelNpu.cpp+1-1

Msrc/aten/src/ATen/native/npu/MseLossBackwardKernelNpu.cpp+2-8

Msrc/aten/src/ATen/native/npu/MseLossKernelNpu.cpp+4-8

Msrc/aten/src/ATen/native/npu/MulKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/MvKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/NeKernelNpu.cpp+3-7

Msrc/aten/src/ATen/native/npu/NonzeroKernelNpu.cpp+1-1

Msrc/aten/src/ATen/native/npu/NormKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/NormalKernelNpu.cpp+4-28

Msrc/aten/src/ATen/native/npu/OneHotKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/OnesLikeKernelNpu.cpp+1-6

Msrc/aten/src/ATen/native/npu/PadKernelNpu.cpp+1-7

Msrc/aten/src/ATen/native/npu/PowKernelNpu.cpp+4-25

Msrc/aten/src/ATen/native/npu/PreluBackwardKernelNpu.cpp+3-9

Msrc/aten/src/ATen/native/npu/PreluKernelNpu.cpp+2-4

Msrc/aten/src/ATen/native/npu/ProdKernelNpu.cpp+8-3

Msrc/aten/src/ATen/native/npu/PtIouKernelNpu.cpp+2-7

Msrc/aten/src/ATen/native/npu/QrKernelNpu.cpp+2-4

Msrc/aten/src/ATen/native/npu/RandomChoiceWithMaskKernelNpu.cpp+3-8

Msrc/aten/src/ATen/native/npu/RandomKernelNpu.cpp+2-5

Msrc/aten/src/ATen/native/npu/RangeKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/ReciprocalKernelNpu.cpp+2-6

Msrc/aten/src/ATen/native/npu/RemainderKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/RepeatKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/RollKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/RoundKernelNpu.cpp+2-11

Msrc/aten/src/ATen/native/npu/RsqrtKernelNpu.cpp+2-10

Msrc/aten/src/ATen/native/npu/RsubKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/ScatterAddKernelNpu.cpp+3-7

Msrc/aten/src/ATen/native/npu/ScatterKernelNpu.cpp+2-6

Msrc/aten/src/ATen/native/npu/ScatterV1KernelNpu.cpp+2-5

Msrc/aten/src/ATen/native/npu/SeluKernelNpu.cpp+1-4

Msrc/aten/src/ATen/native/npu/SigmoidBackwardKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/SigmoidKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/SignKernelNpu.cpp+3-12

Msrc/aten/src/ATen/native/npu/SinhKernelNpu.cpp+1-4

Msrc/aten/src/ATen/native/npu/SliceKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/SmoothL1LossKernelNpu.cpp+3-4

Msrc/aten/src/ATen/native/npu/SoftMarginLossBackwardKernelNpu.cpp+2-8

Msrc/aten/src/ATen/native/npu/SoftmaxBackwardKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/SoftmaxKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/SortKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/SortWithoutIndicesKernelNpu.cpp+2-4

Msrc/aten/src/ATen/native/npu/StackKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/StdKernelNpu.cpp+80-74

Msrc/aten/src/ATen/native/npu/SubKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/SubSampleKernelNpu.cpp+1-4

Msrc/aten/src/ATen/native/npu/SumKernelNpu.cpp+2-1

Msrc/aten/src/ATen/native/npu/TanKernelNpu.cpp+2-12

Msrc/aten/src/ATen/native/npu/TanhBackwardKernelNpu.cpp+2-5

Msrc/aten/src/ATen/native/npu/TanhKernelNpu.cpp+3-13

Msrc/aten/src/ATen/native/npu/ThresholdBackwardKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/ThresholdKernelNpu.cpp+2-10

Msrc/aten/src/ATen/native/npu/TopKKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/TransposeKernelNpu.cpp+2-7

Msrc/aten/src/ATen/native/npu/TrilKernelNpu.cpp+3-10

Msrc/aten/src/ATen/native/npu/TruncKernelNpu.cpp+3-16

Msrc/aten/src/ATen/native/npu/UniformKernelNpu.cpp+2-4

Asrc/aten/src/ATen/native/npu/UpSampleNearest3dKernelNpu.cpp+72-0

Msrc/aten/src/ATen/native/npu/UpsampleBilinear2dKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/UpsampleNearest1dBackwardKernelNpu.cpp+2-4

Msrc/aten/src/ATen/native/npu/UpsampleNearest1dKernelNpu.cpp+2-4

Msrc/aten/src/ATen/native/npu/UpsampleNearest2dBackwardKernelNpu.cpp+2-2

Msrc/aten/src/ATen/native/npu/UpsampleNearest2dKernelNpu.cpp+1-0

Asrc/aten/src/ATen/native/npu/UpsampleNearest3dBackwardKernelNpu.cpp+82-0

Asrc/aten/src/ATen/native/npu/UpsampleTrilinear3dBackwardKernelNpu.cpp+85-0

Asrc/aten/src/ATen/native/npu/UpsampleTrilinear3dKernelNpu.cpp+76-0

Msrc/aten/src/ATen/native/npu/WhereKernelNpu.cpp+2-9

Msrc/aten/src/ATen/native/npu/ZerosLikeKernelNpu.cpp+2-9

Msrc/aten/src/ATen/native/npu/_Unique2KernelNpu.cpp+5-6

Msrc/aten/src/ATen/native/npu/__And__KernelNpu.cpp+4-15

Msrc/aten/src/ATen/native/npu/__Ior__KernelNpu.cpp+1-4

Msrc/aten/src/ATen/native/npu/__Or__KernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/common/MatmulByBmmV2KernelNpu.cpp+13-48

Msrc/aten/src/ATen/native/npu/convolution/Conv3dKernelNpu.cpp+1-8

Msrc/aten/src/ATen/native/npu/frame/OpCommandBase.h+11-5

Msrc/aten/src/ATen/native/npu/frame/OpDynamicParamMaker.h+0-2

Msrc/aten/src/ATen/native/npu/frame/OpParamMaker.cpp+120-4

Msrc/aten/src/ATen/native/npu/frame/OpParamMaker.h+17-109

Msrc/aten/src/ATen/native/npu/interface/EnvVariables.cpp+45-10

Msrc/aten/src/ATen/native/npu/interface/EnvVariables.h+2-0

Rsrc/aten/src/ATen/native/npu/loss/MultilabelMarginLossNpu.cpp→src/aten/src/ATen/native/npu/loss/MultilabelMarginLossKernelNpu.cpp+14-60

Msrc/aten/src/ATen/native/npu/loss/NLLLoss2dKernelNpu.cpp+2-0

Msrc/aten/src/ATen/native/npu/loss/NLLLossBackwardKernelNpu.cpp+10-3

Msrc/aten/src/ATen/native/npu/loss/NLLLossKernelNpu.cpp+11-4

Msrc/aten/src/ATen/native/npu/pooling/AdaptiveAvgPool2dBackwardKernelNpu.cpp+3-11

Msrc/aten/src/ATen/native/npu/pooling/AdaptiveAvgPool3dBackwardKernelNpu.cpp+3-8

Msrc/aten/src/ATen/native/npu/pooling/AvgPool2dKernelNpu.cpp+2-5

Msrc/aten/src/ATen/native/npu/pooling/MaxPool3dWithIndicesBackwardKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/pooling/MaxPool3dWithIndicesKernelNpu.cpp+1-0

Msrc/aten/src/ATen/native/npu/utils/CalcuOpUtil.cpp+2-3

Msrc/aten/src/ATen/native/npu/utils/KernelNpuOutputSize.cpp+2-33

Msrc/aten/src/ATen/native/npu/utils/KernelNpuOutputSize.h+0-9

Rsrc/aten/src/ATen/native/GlobalStep.cpp→src/aten/src/ATen/native/npu/utils/NpuProfilingDispatch.cpp+36-40

Rsrc/aten/src/ATen/native/GlobalStep.h→src/aten/src/ATen/native/npu/utils/NpuProfilingDispatch.h+21-21

Msrc/aten/src/ATen/native/npu/utils/NpuUtils.cpp+1-32

Msrc/aten/src/ATen/native/npu/utils/NpuUtils.h+1-4

Msrc/aten/src/ATen/native/npu/utils/OpPipeWithMultiOut.h+1-2

Msrc/aten/src/ATen/native/npu/utils/OpPreparation.cpp+13-0

Msrc/aten/src/ATen/native/npu/utils/OpPreparation.h+4-0

Msrc/aten/src/ATen/native/npu/utils/OpTemplate.cpp+2-2

Msrc/aten/src/THNPU/THNPUCachingHostAllocator.cpp+10-6

Msrc/c10/npu/NPUCachingAllocator.cpp+8-6

Msrc/c10/npu/NPUEventManager.cpp+8-5

Msrc/c10/npu/NPUStream.cpp+1-1

Msrc/c10/npu/OptionsManager.h+0-1

Msrc/c10/npu/interface/AclInterface.cpp+61-1

Msrc/c10/npu/interface/AclInterface.h+45-2

Msrc/c10/npu/register/FunctionLoader.cpp+0-1

Msrc/third_party/acl/inc/acl/acl.h+0-9

Msrc/third_party/acl/inc/acl/acl_base.h+6-2

Msrc/third_party/acl/libs/acl.cpp+0-2

Msrc/tools/autograd/derivatives.yaml+3-3

Msrc/torch/csrc/autograd/profiler_npu.cpp+6-3

Msrc/torch/npu/__init__.py+2-31

Asrc/torch/npu/npu_frontend_enhance.py+73-0

Mtest/test_npu/test_conv_tbc_backward.py+0-1

Dtest/test_npu/test_instance_norm.py+0-97

Dtest/test_npu/test_instancenorm.py+0-192

Dtest/test_npu/test_multilabel_margin_loss_backward.py+0-176

Mtest/test_npu/test_network_ops/run_tests.py+9-3

Rtest/test_npu/test__Ixor__.py→test/test_npu/test_network_ops/test__Ixor__.py+0-0

Rtest/test_npu/test___iand__.py→test/test_npu/test_network_ops/test___iand__.py+0-1

Rtest/test_npu/test__nnpack_spatial_convolution.py→test/test_npu/test_network_ops/test__nnpack_spatial_convolution.py+0-1

Rtest/test_npu/test_acos.py→test/test_npu/test_network_ops/test_acos.py+0-1

Rtest/test_npu/test_adaptive_avg_pool2d_backward.py→test/test_npu/test_network_ops/test_adaptive_avg_pool2d_backward.py+0-0

Rtest/test_npu/test_adaptive_max_pool2d_backward.py→test/test_npu/test_network_ops/test_adaptive_max_pool2d_backward.py+0-0

Rtest/test_npu/test_addbmm.py→test/test_npu/test_network_ops/test_addbmm.py+0-0

Rtest/test_npu/test_addcdiv.py→test/test_npu/test_network_ops/test_addcdiv.py+0-0

Rtest/test_npu/test_addmv.py→test/test_npu/test_network_ops/test_addmv.py+0-0

Rtest/test_npu/test_addr.py→test/test_npu/test_network_ops/test_addr.py+0-0

Rtest/test_npu/test_affine_grid_generator_backward.py→test/test_npu/test_network_ops/test_affine_grid_generator_backward.py+0-0

Rtest/test_npu/test_asin.py→test/test_npu/test_network_ops/test_asin.py+0-1

Rtest/test_npu/test_bartlett_window.py→test/test_npu/test_network_ops/test_bartlett_window.py+0-1

Rtest/test_npu/test_batch_norm.py→test/test_npu/test_network_ops/test_batch_norm.py+0-0

Rtest/test_npu/test_binary_cross_entropy_with_logits.py→test/test_npu/test_network_ops/test_binary_cross_entropy_with_logits.py+0-0

Mtest/test_npu/test_network_ops/test_bitwise_not.py+0-1

Rtest/test_npu/test_blackman_window.py→test/test_npu/test_network_ops/test_blackman_window.py+0-1

Mtest/test_npu/test_network_ops/test_bmm_v2.py+1-1

Rtest/test_npu/test_cast_Byte.py→test/test_npu/test_network_ops/test_cast_Byte.py+1-1

Rtest/test_npu/test_cast_Char.py→test/test_npu/test_network_ops/test_cast_Char.py+0-0

Rtest/test_npu/test_cast_Float.py→test/test_npu/test_network_ops/test_cast_Float.py+0-0

Rtest/test_npu/test_cast_Half.py→test/test_npu/test_network_ops/test_cast_Half.py+0-0

Rtest/test_npu/test_cast_Int.py→test/test_npu/test_network_ops/test_cast_Int.py+0-0

Rtest/test_npu/test_cast_Long.py→test/test_npu/test_network_ops/test_cast_Long.py+0-0

Rtest/test_npu/test_cast_Short.py→test/test_npu/test_network_ops/test_cast_Short.py+0-0

Rtest/test_npu/test_cdist.py→test/test_npu/test_network_ops/test_cdist.py+0-1

Rtest/test_npu/test_cdist_backward.py→test/test_npu/test_network_ops/test_cdist_backward.py+0-1

Rtest/test_npu/test_celu.py→test/test_npu/test_network_ops/test_celu.py+78-6

Rtest/test_npu/test_conv_tbc.py→test/test_npu/test_network_ops/test_conv_tbc.py+0-7

Rtest/test_npu/test_conv_transpose2d.py→test/test_npu/test_network_ops/test_conv_transpose2d.py+33-21

Rtest/test_npu/test_convolution_backward_weight.py→test/test_npu/test_network_ops/test_convolution_backward_weight.py+3-4

Rtest/test_npu/test_convolution_transpose_backward_weight.py→test/test_npu/test_network_ops/test_convolution_transpose_backward_weight.py+0-1

Rtest/test_npu/test_copy.py→test/test_npu/test_network_ops/test_copy.py+0-1

Rtest/test_npu/test_cos.py→test/test_npu/test_network_ops/test_cos.py+0-1

Rtest/test_npu/test_cosh.py→test/test_npu/test_network_ops/test_cosh.py+0-1

Rtest/test_npu/test_cosinesimilarity.py→test/test_npu/test_network_ops/test_cosinesimilarity.py+0-0

Rtest/test_npu/test_cross.py→test/test_npu/test_network_ops/test_cross.py+0-0

Rtest/test_npu/test_cudnn_convolution_backward_bias.py→test/test_npu/test_network_ops/test_cudnn_convolution_backward_bias.py+0-1

Rtest/test_npu/test_cudnn_convolution_transpose_backward_bias.py→test/test_npu/test_network_ops/test_cudnn_convolution_transpose_backward_bias.py+0-1

Rtest/test_npu/test_cumprod.py→test/test_npu/test_network_ops/test_cumprod.py+0-0

Rtest/test_npu/test_cumsum.py→test/test_npu/test_network_ops/test_cumsum.py+0-0

Rtest/test_npu/test_dim_arange.py→test/test_npu/test_network_ops/test_dim_arange.py+0-0

Rtest/test_npu/test_diml.py→test/test_npu/test_network_ops/test_diml.py+0-0

Rtest/test_npu/test_dirichlet_grad.py→test/test_npu/test_network_ops/test_dirichlet_grad.py+0-0

Rtest/test_npu/test_dot.py→test/test_npu/test_network_ops/test_dot.py+0-1

Rtest/test_npu/test_elu.py→test/test_npu/test_network_ops/test_elu.py+0-0

Rtest/test_npu/test_embedding.py→test/test_npu/test_network_ops/test_embedding.py+0-0

Rtest/test_npu/test_equal.py→test/test_npu/test_network_ops/test_equal.py+0-0

Rtest/test_npu/test_erf.py→test/test_npu/test_network_ops/test_erf.py+0-1

Rtest/test_npu/test_erfinv.py→test/test_npu/test_network_ops/test_erfinv.py+0-1

Rtest/test_npu/test_expm1.py→test/test_npu/test_network_ops/test_expm1.py+0-1

Rtest/test_npu/test_eye.py→test/test_npu/test_network_ops/test_eye.py+0-1

Rtest/test_npu/test_floordivide.py→test/test_npu/test_network_ops/test_floordivide.py+0-0

Rtest/test_npu/test_frac.py→test/test_npu/test_network_ops/test_frac.py+0-1

Rtest/test_npu/test_frobenius_norm.py→test/test_npu/test_network_ops/test_frobenius_norm.py+0-1

Rtest/test_npu/test_full_like.py→test/test_npu/test_network_ops/test_full_like.py+0-1

Rtest/test_npu/test_gelu.py→test/test_npu/test_network_ops/test_gelu.py+0-0

Rtest/test_npu/test_glu.py→test/test_npu/test_network_ops/test_glu.py+0-1

Rtest/test_npu/test_glugrad.py→test/test_npu/test_network_ops/test_glugrad.py+0-1

Rtest/test_npu/test_grid_sampler_2d.py→test/test_npu/test_network_ops/test_grid_sampler_2d.py+0-1

Rtest/test_npu/test_grid_sampler_2d_backward.py→test/test_npu/test_network_ops/test_grid_sampler_2d_backward.py+0-1

Rtest/test_npu/test_group_norm.py→test/test_npu/test_network_ops/test_group_norm.py+0-1

Rtest/test_npu/test_hamming_window.py→test/test_npu/test_network_ops/test_hamming_window.py+0-1

Rtest/test_npu/test_hammingwindow.py→test/test_npu/test_network_ops/test_hammingwindow.py+0-0

Rtest/test_npu/test_hanning_window.py→test/test_npu/test_network_ops/test_hanning_window.py+0-1

Rtest/test_npu/test_hard_sigmoid_backward.py→test/test_npu/test_network_ops/test_hard_sigmoid_backward.py+0-0

Rtest/test_npu/test_hardshrink.py→test/test_npu/test_network_ops/test_hardshrink.py+0-0

Rtest/test_npu/test_hardshrink_backward.py→test/test_npu/test_network_ops/test_hardshrink_backward.py+0-0

Rtest/test_npu/test_hardsigmoid.py→test/test_npu/test_network_ops/test_hardsigmoid.py+0-0

Rtest/test_npu/test_hinge_embedding_loss.py→test/test_npu/test_network_ops/test_hinge_embedding_loss.py+0-0

Rtest/test_npu/test_im2col.py→test/test_npu/test_network_ops/test_im2col.py+0-0

Rtest/test_npu/test_index_fill_d.py→test/test_npu/test_network_ops/test_index_fill_d.py+0-0

Rtest/test_npu/test_index_select.py→test/test_npu/test_network_ops/test_index_select.py+0-0

Rtest/test_npu/test_isclose.py→test/test_npu/test_network_ops/test_isclose.py+0-0

Rtest/test_npu/test_kthvalue.py→test/test_npu/test_network_ops/test_kthvalue.py+0-1

Rtest/test_npu/test_l1_loss.py→test/test_npu/test_network_ops/test_l1_loss.py+0-0

Rtest/test_npu/test_l1_loss_backward.py→test/test_npu/test_network_ops/test_l1_loss_backward.py+0-0

Rtest/test_npu/test_leaky_relu.py→test/test_npu/test_network_ops/test_leaky_relu.py+0-1

Atest/test_npu/test_network_ops/test_log10.py+183-0

Mtest/test_npu/test_network_ops/test_log1p.py+0-1

Rtest/test_npu/test_log_sigmoid.py→test/test_npu/test_network_ops/test_log_sigmoid.py+0-1

Rtest/test_npu/test_log_sigmoid_backward.py→test/test_npu/test_network_ops/test_log_sigmoid_backward.py+0-1

Rtest/test_npu/test_logical_not.py→test/test_npu/test_network_ops/test_logical_not.py+0-1

Rtest/test_npu/test_logsigmoid.py→test/test_npu/test_network_ops/test_logsigmoid.py+0-1

Rtest/test_npu/test_logsigmoidforward.py→test/test_npu/test_network_ops/test_logsigmoidforward.py+0-1

Rtest/test_npu/test_logsumexp.py→test/test_npu/test_network_ops/test_logsumexp.py+0-0

Atest/test_npu/test_network_ops/test_masked_fill_range.py+101-0

Rtest/test_npu/test_max_pool2d_backward.py→test/test_npu/test_network_ops/test_max_pool2d_backward.py+0-0

Rtest/test_npu/test_miopen_batch_norm.py→test/test_npu/test_network_ops/test_miopen_batch_norm.py+0-0

Rtest/test_npu/test_miopen_batch_norm_backward.py→test/test_npu/test_network_ops/test_miopen_batch_norm_backward.py+0-0

Rtest/test_npu/test_miopen_convolution.py→test/test_npu/test_network_ops/test_miopen_convolution.py+0-1

Rtest/test_npu/test_miopen_convolution_backward.py→test/test_npu/test_network_ops/test_miopen_convolution_backward.py+0-1

Rtest/test_npu/test_miopen_convolution_backward_bias.py→test/test_npu/test_network_ops/test_miopen_convolution_backward_bias.py+0-1

Rtest/test_npu/test_miopen_convolution_backward_input.py→test/test_npu/test_network_ops/test_miopen_convolution_backward_input.py+0-1

Rtest/test_npu/test_miopen_convolution_backward_weight.py→test/test_npu/test_network_ops/test_miopen_convolution_backward_weight.py+0-1

Rtest/test_npu/test_mkldnn_adaptive_avg_pool2d.py→test/test_npu/test_network_ops/test_mkldnn_adaptive_avg_pool2d.py+0-1

Rtest/test_npu/test_mkldnn_convolution_backward.py→test/test_npu/test_network_ops/test_mkldnn_convolution_backward.py+0-1

Rtest/test_npu/test_mkldnn_convolution_backward_input.py→test/test_npu/test_network_ops/test_mkldnn_convolution_backward_input.py+0-1

Rtest/test_npu/test_mkldnn_convolution_backward_weights.py→test/test_npu/test_network_ops/test_mkldnn_convolution_backward_weights.py+0-1

Rtest/test_npu/test_multilabel_margin_loss.py→test/test_npu/test_network_ops/test_multilabel_margin_loss.py+0-0

Rtest/test_npu/test_multinomial.py→test/test_npu/test_network_ops/test_multinomial.py+0-0

Rtest/test_npu/test_narrow_copy.py→test/test_npu/test_network_ops/test_narrow_copy.py+0-0

Rtest/test_npu/test_nllloss2d.py→test/test_npu/test_network_ops/test_nllloss2d.py+0-0

Rtest/test_npu/test_norm_except_dim.py→test/test_npu/test_network_ops/test_norm_except_dim.py+1-2

Rtest/test_npu/test_norm_ext.py→test/test_npu/test_network_ops/test_norm_ext.py+3-4

Rtest/test_npu/test_one_hot.py→test/test_npu/test_network_ops/test_one_hot.py+0-1

Rtest/test_npu/test_ones.py→test/test_npu/test_network_ops/test_ones.py+0-0

Rtest/test_npu/test_pdist.py→test/test_npu/test_network_ops/test_pdist.py+0-1

Rtest/test_npu/test_pixel_shuffle.py→test/test_npu/test_network_ops/test_pixel_shuffle.py+0-1

Mtest/test_npu/test_network_ops/test_pooling.py+0-1

Rtest/test_npu/test_prelu.py→test/test_npu/test_network_ops/test_prelu.py+0-1

Rtest/test_npu/test_prelu_backward.py→test/test_npu/test_network_ops/test_prelu_backward.py+0-1

Rtest/test_npu/test_quantize_per_channel.py→test/test_npu/test_network_ops/test_quantize_per_channel.py+0-0

Rtest/test_npu/test_quantize_per_tensor.py→test/test_npu/test_network_ops/test_quantize_per_tensor.py+0-0

Rtest/test_npu/test_real.py→test/test_npu/test_network_ops/test_real.py+0-0

Rtest/test_npu/test_renorm.py→test/test_npu/test_network_ops/test_renorm.py+0-1

Rtest/test_npu/test_repeat_interleave.py→test/test_npu/test_network_ops/test_repeat_interleave.py+0-0

Mtest/test_npu/test_network_ops/test_resize_.py+0-1

Rtest/test_npu/test_roll.py→test/test_npu/test_network_ops/test_roll.py+0-0

Rtest/test_npu/test_selu.py→test/test_npu/test_network_ops/test_selu.py+0-0

Rtest/test_npu/test_sinh.py→test/test_npu/test_network_ops/test_sinh.py+0-0

Rtest/test_npu/test_slow_conv_dilated2d.py→test/test_npu/test_network_ops/test_slow_conv_dilated2d.py+0-0

Rtest/test_npu/test_slow_conv_dilated2d_backward.py→test/test_npu/test_network_ops/test_slow_conv_dilated2d_backward.py+0-0

Rtest/test_npu/test_slow_conv_transpose2d.py→test/test_npu/test_network_ops/test_slow_conv_transpose2d.py+0-0

Rtest/test_npu/test_slow_conv_transpose2d_backward.py→test/test_npu/test_network_ops/test_slow_conv_transpose2d_backward.py+0-0

Rtest/test_npu/test_slow_conv_transpose3d.py→test/test_npu/test_network_ops/test_slow_conv_transpose3d.py+0-1

Rtest/test_npu/test_soft_margin_loss.py→test/test_npu/test_network_ops/test_soft_margin_loss.py+3-4

Rtest/test_npu/test_softmax_backward.py→test/test_npu/test_network_ops/test_softmax_backward.py+0-0

Rtest/test_npu/test_softshrink.py→test/test_npu/test_network_ops/test_softshrink.py+0-1

Rtest/test_npu/test_softshrink_backward.py→test/test_npu/test_network_ops/test_softshrink_backward.py+0-1

Rtest/test_npu/test_split_with_sizes.py→test/test_npu/test_network_ops/test_split_with_sizes.py+0-1

Rtest/test_npu/test_square.py→test/test_npu/test_network_ops/test_square.py+0-0

Mtest/test_npu/test_network_ops/test_std.py+52-1

Rtest/test_npu/test_sum_to_size.py→test/test_npu/test_network_ops/test_sum_to_size.py+0-1

Rtest/test_npu/test_take.py→test/test_npu/test_network_ops/test_take.py+0-0

Rtest/test_npu/test_tan.py→test/test_npu/test_network_ops/test_tan.py+0-0

Rtest/test_npu/test_tensor_npu.py→test/test_npu/test_network_ops/test_tensor_npu.py+8-7

Rtest/test_npu/test_thnn_conv_depthwise2d_backward.py→test/test_npu/test_network_ops/test_thnn_conv_depthwise2d_backward.py+0-0

Rtest/test_npu/test_thnn_conv_depthwise2d_forward.py→test/test_npu/test_network_ops/test_thnn_conv_depthwise2d_forward.py+0-0

Rtest/test_npu/test_threshold_grad_v2_d.py→test/test_npu/test_network_ops/test_threshold_grad_v2_d.py+0-1

Rtest/test_npu/test_threshold_v2_d.py→test/test_npu/test_network_ops/test_threshold_v2_d.py+0-0

Rtest/test_npu/test_trapz_dx.py→test/test_npu/test_network_ops/test_trapz_dx.py+1-2

Rtest/test_npu/test_trapz_x.py→test/test_npu/test_network_ops/test_trapz_x.py+0-1

Rtest/test_npu/test_true_divide.py→test/test_npu/test_network_ops/test_true_divide.py+0-1

Rtest/test_npu/test_trunc.py→test/test_npu/test_network_ops/test_trunc.py+0-1

Mtest/test_npu/test_network_ops/test_unbind_copy_contiguous.py+0-1

Rtest/test_npu/test_upsample_bilinear2d.py→test/test_npu/test_network_ops/test_upsample_bilinear2d.py+0-1

Rtest/test_npu/test_upsample_linear1d.py→test/test_npu/test_network_ops/test_upsample_linear1d.py+0-1

Rtest/test_npu/test_upsample_nearest2d.py→test/test_npu/test_network_ops/test_upsample_nearest2d.py+0-0

Rtest/test_npu/test_upsample_nearest2d_backward.py→test/test_npu/test_network_ops/test_upsample_nearest2d_backward.py+0-0

Atest/test_npu/test_network_ops/test_upsample_nearest3d.py+82-0

Atest/test_npu/test_network_ops/test_upsample_nearest3d_backward.py+92-0

Atest/test_npu/test_network_ops/test_upsample_trilinear3d.py+82-0

Atest/test_npu/test_network_ops/test_upsample_trilinear3d_backward.py+91-0

Dtest/test_npu/test_trace.py+0-87