[RFC]: 通过自定义方式接入Triton算子，支持静态图

# 背景与目标描述. 当前Triton算子生态较好，昇腾团队也提供了[triton_ascend插件](https://gitcode.com/Ascend/triton-ascend)用于支持昇腾算子。但该插件仅支持通过torch_npu接入，我们希望在mindspore上也能通过自定义的方式接入triton-ascend来复用triton算子。 # 建议的方案. ## 1. 原生Triton算子调用流程分析 开发Triton算子采用的是纯Python语言，通过`@triton.jit`装饰器包裹triton DSL，在Python调用该接口时会自动触发算子编译，例如： ``` import triton import triton.language as tl @triton.jit def add_kernel(x_ptr, # *Pointer* to first input vector. y_ptr, # *Pointer* to second input vector. output_ptr, # *Pointer* to output vector. n_elements, # Size of the vector. BLOCK_SIZE: tl.constexpr, # Number of elements each program should process. ): # There are multiple 'programs' processing different data. We identify which program # we are here: pid = tl.program_id(axis=0) # We use a 1D launch grid so axis is 0. # This program will process inputs that are offset from the initial data. # For instance, if you had a vector of length 256 and block_size of 64, the programs # would each access the elements [0:64, 64:128, 128:192, 192:256]. # Note that offsets is a list of pointers: block_start = pid * BLOCK_SIZE offsets = block_start + tl.arange(0, BLOCK_SIZE) # Create a mask to guard memory operations against out-of-bounds accesses. mask = offsets < n_elements # Load x and y from DRAM, masking out any extra elements in case the input is not a # multiple of the block size. x = tl.load(x_ptr + offsets, mask=mask) y = tl.load(y_ptr + offsets, mask=mask) output = x + y # Write x + y back to DRAM. tl.store(output_ptr + offsets, output, mask=mask) def add(x, y, out): n_elements = out.numel() def grid(meta): return (triton.cdiv(n_elements, meta['BLOCK_SIZE']), ) return add_kernel[grid](x, y, out, n_elements, BLOCK_SIZE=1024) ``` 这里按照torch自定义算子的一般写法，在外部调用add接口时传入的是输入和输出tensor。输出tensor一般是在前面通过`empty`或`empty_like`创建。 triton算子是动态编译的，每次调用时需传入 **真实Tensor** 。算子一般是基于shape来做缓…