docs: add English translations for markdown files-ops-ras-AtomGit

A.gitcode/PULL_REQUEST_TEMPLATE.en-US.md+20-0

@@ -0,0 +1,20 @@
	1	+## Description
	2	+<!--Describe your changes in detail here, including the reason for the changes and the method adopted.-->
	3	+
	4	+## Related Issue
	5	+<!--If this PR is intended to resolve a specific issue, provide the issue link here. For example: Related Issue #000-->
	6	+
	7	+## Tests
	8	+<!--Describe the tests performed to verify your changes.-->
	9	+
	10	+## Documentation Update
	11	+<!--If this PR contains documentation updates, specify them here. For example: the README.md file is updated.-->
	12	+
	13	+## Type Label
	14	+<!-- [x] means selected -->
	15	+- [ ] Bug fix
	16	+- [ ] New feature
	17	+- [ ] Performance optimization
	18	+- [ ] Documentation update
	19	+- [ ] Other, describe:
	20	+

Aexamples/README_en.md+36-0

@@ -0,0 +1,36 @@
	1	+English \| [简体中文](./README.md)
	2	+
	3	+## Overview
	4	+
	5	+> Note:
	6	+> - Currently, most operators in the operator library run on AI Core, and a few run on AI CPU. By default, an operator mentioned in this project refers to an AI Core operator.
	7	+> - For the introduction to AI Core and AI CPU, see the "Concepts, Principles, and Terms > Hardware Architecture and Data Processing Principles" section in [Ascend C Operator Development](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC).
	8	+
	9	+This project provides development and invocation samples for AI Core operators and AI CPU operators. Refer to the corresponding implementations based on actual requirements.
	10	+
	11	+## Directory Description
	12	+```
	13	+├── examples
	14	+│ ├── add_example # AI Core operator name
	15	+│ │ ├── CMakeLists.txt # Operator compilation configuration file; keep the original file
	16	+│ │ ├── examples # Operator usage samples
	17	+│ │ ├── op_graph # Directory related to operator graph construction
	18	+│ │ ├── op_host # Implementation of the operator information library, Tiling, and InferShape
	19	+│ │ └── op_kernel # Operator kernel directory
	20	+│ ├── add_example_aicpu # AI CPU operator name
	21	+│ │ ├── CMakeLists.txt # Operator compilation configuration file; keep the original file
	22	+│ │ ├── examples # Operator usage samples
	23	+│ │ ├── op_graph # Directory related to operator graph construction
	24	+│ │ ├── op_host # Implementation of the operator information library and InferShape
	25	+│ │ └── op_kernel_aicpu # Operator kernel directory
	26	+│ ├── CMakeLists.txt # Operator compilation configuration file; keep the original file
	27	+│ └── README.md # Operator description document
	28	+
	29	+```
	30	+
	31	+## Operator Development Samples
	32	+\|Sample Directory\|Sample Description\|Operator Development\|Operator Invocation\|
	33	+\|---\|------------------\|---\|---\|
	34	+\| add_example \| Operator that implements the addition of two tensors. \| For the end-to-end operator development process, see [AI Core Operator Development Guide](../docs/en/develop/aicore_develop_guide.md). \| For invocation samples, see [examples](./add_example/examples/). \|
	35	+\|add_example_aicpu\| Operator that implements the addition of two tensors. \| For the end-to-end operator development process, see [AI CPU Operator Development Guide](../docs/en/develop/aicpu_develop_guide.md).\| For invocation samples, see [examples](./add_example_aicpu/examples/). \|
	36	+\|fast_kernel_launch_example\| Sample that implements fast end-to-end operator development in PyTorch scenarios. \| For the end-to-end operator development process, see [PyTorch Operator Quick Development Guide](./fast_kernel_launch_example/README_en.md).\| For invocation samples, see [examples](./fast_kernel_launch_example/). \|

Mexamples/add_example/README_en.md+2-0

	@@ -1,5 +1,7 @@
1	# AddExample	1	# AddExample
2		2
		3	+English \| [简体中文](./README.md)
		4	+
3	## Product Support	5	## Product Support
4		6
5	\| Product \| Supported \|	7	\| Product \| Supported \|

Mexamples/add_example_aicpu/README_en.md+2-0

	@@ -1,5 +1,7 @@
1	# AddExample	1	# AddExample
2		2
		3	+English \| [简体中文](./README.md)
		4	+
3	## Product Support	5	## Product Support
4		6
5	\| Product \| Supported \|	7	\| Product \| Supported \|

Aexamples/fast_kernel_launch_example/README_en.md+190-0

@@ -0,0 +1,190 @@
	1	+# AscendOps
	2	+
	3	+English \| [简体中文](./README.md)
	4	+
	5	+## Prerequisites
	6	+
	7	+- Before you start, complete the basic environment setup by referring to [Environment Setup](../../docs/en/install/quick_install.md).
	8	+- GCC 9.4.0+
	9	+- Python 3.8+
	10	+- PyTorch>=2.6.0
	11	+- [TorchNPU](https://gitcode.com/Ascend/pytorch/releases) of the matching version
	12	+
	13	+## Installation Steps
	14	+
	15	+1. Install Dependencies:
	16	+ ```sh
	17	+ python3 -m pip install -r requirements.txt
	18	+ ```
	19	+
	20	+2. Build the Wheel:
	21	+ ```sh
	22	+ # -n: non-isolated build (uses existing environment)
	23	+ python3 -m build --wheel -n
	24	+ ```
	25	+
	26	+3. Install Package:
	27	+ ```sh
	28	+ python3 -m pip install dist/*.whl --force-reinstall --no-deps
	29	+ ```
	30	+
	31	+4. (Optional) Before rebuilding, it is recommended that you run the following command to clear the compilation cache:
	32	+ ```sh
	33	+ python3 setup.py clean
	34	+ ```
	35	+
	36	+## Quick Start
	37	+After the installation, you can use NPU operators in the same way as ordinary PyTorch operations.
	38	+
	39	+```python
	40	+import torch
	41	+import torch_npu
	42	+import ascend_ops
	43	+
	44	+# Initialize data on NPU
	45	+x = torch.randn(10, 32, dtype=torch.float32).npu()
	46	+y = torch.randn(10, 32, dtype=torch.float32).npu()
	47	+
	48	+# Call the custom NPU operator
	49	+npu_result = torch.ops.ascend_ops.add(x, y)
	50	+
	51	+# Verify against CPU ATen implementation
	52	+cpu_x = x.cpu()
	53	+cpu_y = y.cpu()
	54	+cpu_result = cpu_x + cpu_y
	55	+
	56	+assert torch.allclose(cpu_result, npu_result.cpu(), rtol=1e-6)
	57	+print("Verification successful!")
	58	+```
	59	+
	60	+## Developer Guide: Adding a New Operator
	61	+
	62	+To implement a new operator (for example, `add`), you only need to provide a C++ implementation.
	63	+
	64	+1. First, create a folder named after the operator `add` in the csrc directory. In this folder, create a subfolder `ascend910b` named after the SoC for which you want to develop.
	65	+
	66	+2. Create a `CMakeLists.txt` file in the soc directory:
	67	+ ```
	68	+ add_sources("--npu-arch=dav-2201")
	69	+ ```
	70	+ Here, `dav-2201` is the compilation parameter corresponding to the ascend910b chip.
	71	+
	72	+3. Create an `add.cpp` file in the soc directory (using the operator name as the file name is recommended). This file contains all modules required for developing an AI Core operator.
	73	+ - Operator schema registration
	74	+ - Operator Meta Function implementation & registration
	75	+ - Operator kernel implementation (AscendC)
	76	+ - Operator NPU invocation implementation & registration
	77	+
	78	+ ```cpp
	79	+ #include <ATen/Operators.h>
	80	+ #include <torch/all.h>
	81	+ #include <torch/library.h>
	82	+ #include "torch_npu/csrc/core/npu/NPUStream.h"
	83	+ #include "torch_npu/csrc/framework/OpCommand.h"
	84	+ #include "kernel_operator.h"
	85	+ #include "platform/platform_ascendc.h"
	86	+ #include <type_traits>
	87	+
	88	+ namespace ascend_ops { // The current project is a namespace
	89	+ namespace Add { // Each operator is recommended to have its own namespace to prevent global variable pollution
	90	+
	91	+ /**
	92	+ * Register the operator schema with the PyTorch framework
	93	+ * so that the framework knows that this operator exists
	94	+ */
	95	+ // Register the operator's schema
	96	+ TORCH_LIBRARY_FRAGMENT(EXTENSION_MODULE_NAME, m)
	97	+ {
	98	+ m.def("add(Tensor x, Tensor y) -> Tensor");
	99	+ }
	100	+
	101	+ /**
	102	+ * Implement the Meta function of the operator, that is, InferShape+InferDtype
	103	+ * Infer the output and the required space of the operator from the inputs, without actually computing the operator
	104	+ */
	105	+ // Meta function implementation of Add
	106	+ torch::Tensor add_meta(const torch::Tensor &x, const torch::Tensor &y)
	107	+ {
	108	+ TORCH_CHECK(x.sizes() == y.sizes(), "The shapes of x and y must be the same.");
	109	+ auto z = torch::empty_like(x);
	110	+ return z;
	111	+ }
	112	+
	113	+ /**
	114	+ * Register the Meta function of the operator with the framework
	115	+ * The framework can call this Meta function to know the required space before actually executing the operator computation
	116	+ * This lays the foundation for graph acceleration such as torch.compile/AutoGrad/AclGraph
	117	+ */
	118	+ // Register the Meta implementation
	119	+ TORCH_LIBRARY_IMPL(EXTENSION_MODULE_NAME, Meta, m)
	120	+ {
	121	+ m.impl("add", add_meta);
	122	+ }
	123	+
	124	+ /**
	125	+ * NPU operator kernel implementation, written in AscendC APIs for the current soc
	126	+ */
	127	+ template <typename T>
	128	+ __global__ __aicore__ void add_kernel(GM_ADDR x, GM_ADDR y, GM_ADDR z, int64_t totalLength, int64_t blockLength, uint32_t tileSize)
	129	+ {
	130	+ // kernel implementation
	131	+ }
	132	+
	133	+ /**
	134	+ * Implement the operator invocation interface
	135	+ * In this interface, the NPU kernel invocation needs to be completed
	136	+ * 1. Compute the number/shape/dtype of the output tensors (by calling the Meta function or implementing directly)
	137	+ * 2. Compute Tiling: determine how to split the computation into blocks based on the shape
	138	+ * 3. Invoke the NPU kernel
	139	+ *
	140	+ */
	141	+ torch::Tensor add_npu(const torch::Tensor &x, const torch::Tensor &y)
	142	+ {
	143	+ auto z = add_meta(x, y);
	144	+ auto stream = c10_npu::getCurrentNPUStream().stream(false);
	145	+ int64_t totalLength, blockDim, blockLength, tileSize;
	146	+ totalLength = x.numel();
	147	+ std::tie(blockDim, blockLength, tileSize) = calc_tiling_params(totalLength);
	148	+ auto x_ptr = (GM_ADDR)x.data_ptr();
	149	+ auto y_ptr = (GM_ADDR)y.data_ptr();
	150	+ auto z_ptr = (GM_ADDR)z.data_ptr();
	151	+ auto acl_call = [=]() -> int {
	152	+ AT_DISPATCH_SWITCH(
	153	+ x.scalar_type(), "add_npu",
	154	+ // Invoke different NPU kernels for different data types
	155	+ AT_DISPATCH_CASE(torch::kFloat32, [&] {
	156	+ using scalar_t = float;
	157	+ add_kernel<scalar_t><<<blockDim, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);
	158	+ })
	159	+ AT_DISPATCH_CASE(torch::kFloat16, [&] {
	160	+ using scalar_t = half;
	161	+ add_kernel<scalar_t><<<blockDim, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);
	162	+ })
	163	+ AT_DISPATCH_CASE(torch::kInt32, [&] {
	164	+ using scalar_t = int32_t;
	165	+ add_kernel<scalar_t><<<blockDim, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);
	166	+ })
	167	+ );
	168	+ return 0;
	169	+ };
	170	+ // RunOpApi/RunOpApiV2 must be used for the invocation to ensure that the timing is consistent with that of aclnn API calls made by TorchNPU.
	171	+ at_npu::native::OpCommand::RunOpApi("Add", acl_call);
	172	+ return z;
	173	+ }
	174	+
	175	+ /**
	176	+ * Register the invocation function of the operator with the framework, where Device is PrivateUse1
	177	+ * The framework knows that when all inputs are on the NPU device, it dispatches to this operator implementation
	178	+ */
	179	+ // Register the NPU implementation
	180	+ TORCH_LIBRARY_IMPL(EXTENSION_MODULE_NAME, PrivateUse1, m)
	181	+ {
	182	+ m.impl("add", add_npu);
	183	+ }
	184	+
	185	+ } // namespace Add
	186	+ } // namespace ascend_ops
	187	+
	188	+ ```
	189	+4. Build the wheel package by referring to the [Installation Steps](#installation-steps) section, and then install and test it.
	190	+5. For testing the operator API, refer to the implementation of [test_add.py](tests/add/test_add.py)