做分类信息网站成功的品牌推广案例分析

做分类信息网站,成功的品牌推广案例分析,wordpress伪静态iis7,品牌网站设计网站00 概要Mirage 使用 uGraph 来指定在 GPU 上执行张量程序。uGraph 包含多个级别的层次化图#xff0c;以表示在内核、块和线程级别的计算。下图是GQA对应的μGraphs#xff0c;显示了一个用于计算GQA的 uGraph。我们用它作为运行示例来解释 uGraph 的关键组成部分。mugraph_g…00 概要Mirage 使用 uGraph 来指定在 GPU 上执行张量程序。uGraph 包含多个级别的层次化图以表示在内核、块和线程级别的计算。下图是GQA对应的μGraphs显示了一个用于计算GQA的 uGraph。我们用它作为运行示例来解释 uGraph 的关键组成部分。mugraph_gqa0x01 机制1.1 当前问题LLM 的计算过程通常以计算图的形式表示其中每个节点对应一个计算算子如矩阵乘法、注意力机制或集合通信原语如 all-reduce边表示算子间的数据依赖关系。现有系统通常为每个算子启动独立的 GPU 内核。然而这种“单算子单内核”的执行模型难以实现 pipeline 优化因为依赖关系是在整个内核的粗粒度层面强制执行的而非实际数据单元层面。例如矩阵乘法matmul后接 all-reduce 操作现有系统中all-reduce 内核必须等待整个 matmul 内核完成。而实际上all-reduce 的每个数据分块仅依赖 matmul 输出的局部结果。这种逻辑依赖与实际依赖的错配严重限制了计算与通信的重叠潜力。下图的右侧展示次优方案 —— 其引入不必要的数据依赖与全局屏障导致跨层流水线优化机会受限。img1.2 解决方案为了解决这一问题Mirage实现了多层次计算图表示μGraphs与归纳式程序合成Inductive Program Synthesis。这两大机制协同作用实现了从宏观调度到微观计算的全链路优化高效生成GPU程序显著提升了张量计算的性能。Mirage 的编译流程清晰且目标明确输入来自预定义算子集合的计算图子图如 GQA 注意力计算子图确保输入逻辑的规范性与可优化性核心优化步骤包含图重写Graph Rewrite调整图结构以适配 GPU 架构、算子融合Operator Fusion减少内存访问次数等所有优化均基于 μGraphs 的跨层级表示展开输出优化后的 CUDA 程序直接适配 GPU 硬件执行可直接JIT嵌入pytorch。1.2.1 μGraphs多层次计算图表示MPK 编译器将 LLM 计算图自动转化为细粒度任务图最大化暴露并行性。该任务图在子内核级别显式捕获依赖关系实现更激进的跨层流水线优化。具体而言在 MPK 任务图中参见上图任务矩形表示代表分配给单个 GPU 流式多处理器SM的计算或通信单元。事件圆形表示表示任务间的同步点。触发机制每个任务发出指向触发事件的边该事件在关联任务全部完成后激活。依赖机制每个任务接收来自依赖事件的边表明事件激活后任务立即启动。任务图使 MPK 能够发掘计算图中无法实现的 pipeline 优化机会。例如MPK 可以构建优化任务图 —— 其中每个 all-reduce 任务仅依赖于生成其输入的对应 matmul 任务从而实现分块执行与计算通信重叠。除生成优化任务图外MPK 还通过 Mirage 内核超优化器自动为每个任务生成高性能 CUDA 实现确保任务在 GPU 流式多处理器SM上高效执行。1.2.2 归纳式程序合成优化范式归纳式程序合成是Mirage的另一大核心机制。与传统的演绎式程序合成如基于规则的重写系统不同归纳式程序合成直接从语法出发构造程序并借助SMT求解器验证构造程序与原程序的等价性。这种方法能够突破传统优化方法的局限发现将代数变换、调度变换和新自定义内核生成相结合的创新优化路径。通过归纳式程序合成Mirage能够自动生成高性能的GPU内核代码不仅简化了开发流程还提升了程序的运行效率使得开发者能够更专注于高层逻辑的设计而无需深入底层硬件细节。传统机器学习编译器如 TVM、TensorRT采用演绎式程序合成Deductive Program Synthesis又称 Term Rewrite 从原始程序出发通过等价重写规则如图模式匹配、循环调度原语逐步变换始终在 “程序等价类” 内搜索更优实现 —— 这种方式依赖手工设计规则难以突破现有等价类的性能上限。Mirage 则采用归纳式程序合成不依赖原始程序的逐步变换而是直接基于算子语法构造全新候选程序再通过 “μGraphs 语义校验 概率等价验证”如有限域随机测试确认候选程序与原始程序的功能一致性。这种范式无需受限于等价重写规则可探索更灵活的跨层级优化方案如 Kernel-Graph 合成算子与 Block-Graph 共享内存复用的协同同时通过概率验证保障正确性。下图是Mirage找出的最佳μGraphs。34--PersistentKernel 图0x02 多层次计算图表示Mirage 实现了多层次计算图表示μGraphs通过 kernel-graph内核图、block-graph块图和 thread-graph线程图这三层结构化图模型精确映射 GPU 程序从内核到线程的执行逻辑与存储层级。这种三层结构与 CUDA 程序的执行层级及 GPU 的存储体系紧密对应每层均清晰定义了 “算子类型 — 张量存储 — 核心功能” 的关联关系。2.1 概念三层的概念如下kernel-graph内核图属于高层次抽象用于表示整个计算图即完整的计算任务包含粗粒度的高层操作如完整的矩阵乘法、规约运算等与对应数据。该层负责全局调度重点关注数据流与任务间的依赖关系对应 GPU 的全局内存主要处理宏观层面的任务分配与协同。其包含的算子举例类型有高层操作KN_INPUT_OP输入算子、KN_OUTPUT_OP输出算子、KN_MATMUL_OP矩阵乘法算子数学操作KN_EXP_OP指数运算算子、KN_ADD_OP加法算子、KN_MUL_OP乘法算子规约操作KN_REDUCTION_0_OP零阶规约算子等自定义操作KN_CUSTOMIZED_OP自定义算子等。block-graph块图属于中等层次抽象嵌套在 KN_CUSTOMIZED_OP自定义内核算子中定义 threadblock线程块级别的计算逻辑。该层包含细粒度操作负责管理线程块级别的并行计算重点关注内存访问模式、循环结构等中观细节对应 GPU 的共享内存核心目标是优化中观层面的资源利用与数据共享效率。其包含的算子类型举例有输入操作TB_INPUT_OP线程块输入算子内存操作TB_MATMUL_OP线程块矩阵乘法算子、TB_EXP_OP线程块指数运算算子特殊操作TB_FORLOOP_ACCUM_NO_RED_OP线程块循环累加无规约算子、TB_RMS_NORM_OP线程块 RMS 归一化算子。thread-graph线程图在 block-graph 的具体操作中体现定义线程级别的执行细节。该层专注于线程级别的微观计算逻辑对应 GPU 的寄存器核心作用是确保每个线程的高效执行最大化单线程的计算吞吐量。这种三层结构支持系统在不同抽象层级开展针对性优化在 kernel-graph 层主要进行全局任务调度与数据流优化明确整体计算流程与资源分配方向在 block-graph 层侧重线程块级别的并行策略优化提升中观层面的并行效率与数据共享能力在 thread-graph 层聚焦具体的内存访问模式优化与计算指令调度确保微观执行的高效性。若用通俗语言概括三层结构的分工kernel-graph 决定 “要做什么”明确整体计算任务与目标block-graph 决定 “该怎么做”规划线程块级的执行方案thread-graph 负责 “具体执行”完成线程级的微观计算。这种从宏观到微观的层次化设计使 μGraphs 能够实现从全局调度到局部执行的全链路优化有效减少计算冗余与资源浪费确保 GPU 计算资源的高效利用。2.2 层级关系三级图结构的关系如下图所示。muGraph(Kernel Graph)│├────► KNOperator(各种标准操作)││└────► KNCustomizeOp(自定义操作)│└───► block-graph(Threadblock Graph)│├────► TBOperator(各种线程块操作)│└────► TBInputOp(连接到muGraph的张量)│└───► thread-level execution(线程级执行)2.3 对比三层的对比如下。计算图层级 对应 CUDA 执行层级 张量存储位置 算子类型与功能 核心属性 / 逻辑Kernel-Graph 整个 GPU 内核多流处理器 SM 协同 设备全局内存Device DRAM 1. 预定义算子直接调用厂商库内核如 cuBLAS 的 GEMM 矩阵乘、cuDNN 的卷积 2. 合成算子需通过更低层级的 Block-Graph 描述承载算子融合、自定义算法等复杂逻辑 无额外属性核心是 “调度多 SM 协同”通过预定义算子复用成熟库性能合成算子支持灵活优化Block-Graph 单个流处理器 SM线程块协作 共享内存Shared Memory 1. 预定义算子调用 CUTLASS、ThunderKittens 等库的共享内存操作如块内矩阵乘、累加 2. 合成算子由 Thread-Graph 描述实现线程块内细粒度计算 1. 并行切分属性imap输入分块映射 Grid 维度到输入张量维度、omap输出拼接映射 Grid 维度到输出张量维度、fmap循环迭代映射 For-Loop 维度到数据迭代器 / 累加器维度 2. 执行逻辑支持线程块循环迭代通过共享内存复用与 “计算 - 访存重叠”将全局内存读写延迟隐藏在计算过程中Thread-Graph 单个线程寄存器操作 线程私有寄存器Register File 仅含预定义算子描述单个线程内的寄存器级流水操作如 load 数据→元素级计算→store 结果支持循环迭代与寄存器累加默认通过 “规则化融合” 快速生成避免细粒度层级的冗余搜索 核心是 “单线程高效流水”通过寄存器操作最小化内存访问提升计算密度2.4 执行关系persistent_kernel.py是 Persistent Kernel的Python接口本质是Python到CUDA持久化内核系统的桥梁允许用户用python定义复杂的计算图然后在GPU上高效执行。persistent_kernel.py与三层计算图的关系如下Persistent Kernel 创建并管理 Kernel GraphKernel Graph 通过 KN_CUSTOMIZED_OP 包含多个 Block Graph每个 Block Graph 定义线程块内的操作序列Kernel Graph 转换为 Task Graph 用于执行Task Execution Engine 在 Persistent Kernel 中执行任务Event System 管理任务间的依赖和同步Thread Graph 在实际GPU线程中执行具体操作0x03 内核图每个张量程序对应一个内核图其中每个节点代表在整個 GPU 上运行的内核每条边是内核之间共享的张量。内核图中的所有张量都存储在 GPU 设备内存中因为不同的内核不能在寄存器文件或共享内存中共享数据。内核图中的每个节点都可以是现有内核库如 cuDNN 的卷积和 cuBLAS 的矩阵乘法支持的预定义内核操作符。此外为了启用细粒度的内核间优化如内核融合内核图中的节点也可以是图定义的内核操作符其语义和行为由较低级别的即块图定义。下图中的两个内核操作符都是图定义的操作符每个都由块图指定。mugraph_gqa3.1 PersistentKernel调用在PersistentKernel内部kn_graph负责实际的计算图构建。self.kn_graph KNGraph(CyKNGraph(disable_fingerprintTrue))每个attach_input和new_tensor调用都会在kn_graph中创建张量节点。每个layer调用也会在kn_graph中添加相应的计算节点。最后compile()调用self.kn_graph.generate_task_graph生成任务图。3.2 Python 代码内核图在Python中的类是KNGraph。KNGraph用于构建和管理内核计算图。比如new_input会创建新的输入变量。attach_torch_tensor管理PyTorch变量。attach_cuda_tensor关联CUDA变量。compile会生成最终的执行代码。KNGraph的特点如下Kernel graph的节点是预定义算子pre-defined operator比如cuBLAS GEMM、cuDNN Conv合成算子graph-defined operator用更低一层的block graph描述可承载fusion/新算法。Kernel graph的边是位于全局内存Device DRAM的Tensor。KNGraph 代码举例如下class KNGraph:def __init__(self, graph):self.cygraph graphself._is_compiled Falseself.run Noneself._valid_cuda_kernels Falseself._cached_results Noneself.visualizer Noneself.backend cudadef new_input(self, dims: tuple, strides: tuple None, dtype: dtype float16) - DTensor:# use the default strided layout if strides Noneif strides is None:total_elements 1strides []for d in reversed(dims):strides.append(total_elements)total_elements * dstrides reversed(strides)return self.cygraph.new_input(dims, tuple(strides), dtype)def compile(self, async_False, **kwargs):if self._is_compiled:return self._cached_resultsinput_tensors kwargs.get(inputs, [])input_strides []for i in range(len(dtensors)):dims, strides self.cygraph.get_input_dtensor_shape_and_stride(dtensors[i])input_strides.append(strides)target_cc kwargs.get(target_cc,torch.cuda.get_device_properties(0).major * 10 torch.cuda.get_device_properties(0).minor,)num_warp_groups kwargs.get(num_warp_groups, 2)pipeline_stages kwargs.get(pipeline_stages, 2)enable_online_softmax kwargs.get(enable_online_softmax, False)result generate_cuda_program(self.cygraph,target_cctarget_cc,input_stridesinput_strides,num_warp_groupsnum_warp_groups,pipeline_stagespipeline_stages,profilingprofiling,enable_online_softmaxenable_online_softmax,)if result[max_smem_size] get_shared_memory_capacity(target_cc):self._is_compiled Trueself._valid_cuda_kernels Falseself._error_message shared memory usage exceed limitif async_:return Handle([], None)else:return NoneMIRAGE_ROOT, INCLUDE_PATH, DEPS_PATH get_key_paths()tempdir_obj tempfile.TemporaryDirectory()tempdir tempdir_obj.namesaved_addr file_id kwargs.get(file_id, -1)if file_id ! -1:print(ffile_id: {file_id})saved_addr f./generated_codes/{file_id}/FILE_NAME os.path.join(tempdir, test.cu)so_path os.path.join(tempdir, test.cpython-38-x86_64-linux-gnu.so)with open(FILE_NAME, w) as f:f.write(result[code] HARD_CODE)if saved_addr ! :print(fsaved_addr: {saved_addr})os.makedirs(saved_addr, exist_okTrue)with open(saved_addr test str(file_id) .cu, w) as f:f.write(result[code] HARD_CODE)cc shutil.which(nvcc)# This function was renamed and made public in Python 3.10if hasattr(sysconfig, get_default_scheme):scheme sysconfig.get_default_scheme()else:scheme sysconfig._get_default_scheme()if scheme posix_local:scheme posix_prefixpy_include_dir sysconfig.get_paths(schemescheme)[include]cc_cmd get_cc_cmd(target_cc,cc,FILE_NAME,py_include_dir,INCLUDE_PATH,DEPS_PATH,so_path,profiling,)def remain_op():import importlib.utiltry:spec importlib.util.spec_from_file_location(__mirage_launcher, so_path)mod importlib.util.module_from_spec(spec)spec.loader.exec_module(mod)self.run getattr(mod, launch)self._is_compiled Trueself._valid_cuda_kernels Trueself._cached_results resultself._error_message No errortempdir_obj.cleanup()return self._cached_resultsexcept ImportError:self._is_compiled Trueself._valid_cuda_kernels Falseself._cached_results Noneself._error_message CUDA compilation errorreturn Noneif async_:if global_config.bypass_compile_errors:ret subprocess.Popen(cc_cmd, stdoutsubprocess.DEVNULL, stderrsubprocess.STDOUT)else:ret subprocess.Popen(cc_cmd)return Handle([ret], remain_op)else:ret subprocess.check_call(cc_cmd)return remain_op()3.3 桥梁PersistentKernel 中通过如下方式进行设置 Kernel Graph。self.kn_graph KNGraph(CyKNGraph(disable_fingerprintTrue))在python\mirage_cython\core.pyx 文件中CyKNGraph 中有定义 CppKNGraph。cdef class CyKNGraph:cdef CppKNGraph *p_kgraph #Hold a CppKNGraph instancedef __cinit__(self, graph None, bool disable_fingerprint False):cdef unsigned long long ptrcdef dim3 c_gpu_dimif graph is None:c_gpu_dim.x 1c_gpu_dim.y 1c_gpu_dim.z 1self.p_kgraph new CppKNGraph(c_gpu_dim, disable_fingerprint)else:ptr ctypes.cast(graph, ctypes.c_void_p).valueself.p_kgraph CppKNGraph*(ptr)在 python\mirage_cython\CCore.pxd 文件中指明 CppKNGraph 对应了 mirage::kernel::Graph这便是C代码中Kernel Graph 的实现。cdef cppclass CppKNGraph mirage::kernel::Graph:CppKNGraph(dim3 gpu_dim, bool disable_fingerprint)CppDTensor* new_input_ptr(vector[int] dims,vector[size_t] strides,DataType data_type,DmemLayout layout)void mark_output(const CppDTensor* A, vector[size_t] strides)CppDTensor* matmul(const CppDTensor* A, const CppDTensor* B)CppDTensor* reduction(const CppDTensor* input, int dim, int size)CppDTensor* rms_norm(const CppDTensor* input, vector[int])CppDTensor* exp(const CppDTensor* input)CppDTensor* silu(const CppDTensor* input)CppDTensor* gelu(const CppDTensor* input)CppDTensor* relu(const CppDTensor* input)CppDTensor* clamp(const CppDTensor* input, float min_val, float max_val)CppDTensor* sqrt(const CppDTensor* input)CppDTensor* square(const CppDTensor* input)CppDTensor* add(const CppDTensor* op1, const CppDTensor* op2)CppDTensor* mul(const CppDTensor* op1, const CppDTensor* op2)CppDTensor* div(const CppDTensor* op1, const CppDTensor* op2)CppDTensor* pow(const CppDTensor* op1, const CppDTensor* op2)int customized(vector[const CppDTensor*] inputs,CppDTensor** outputs,CppTBGraph* bgraph)int get_num_input_dtensors()int get_num_output_dtensors()int get_input_dtensors(CppDTensor** cinputs)int get_input_dtensor_shape_and_stride(const CppDTensor *input, int *strides, int *dims)void generate_triton_program(const char *filepath)void generate_cuda_program(const char *filepath)size_t get_owner_independent_hash() const# Persistent kernel functionsvoid attach_torch_tensor(const CppDTensor *input,void *torch_data_ptr,const char *name)void attach_cuda_tensor(const CppDTensor *input,const char *name)void attach_nvshmem_tensor(const CppDTensor *input,const char *name)CppDTensor* fuse_tensors(vector[const CppDTensor*] inputs,int fused_dim,int num_groups,const char *name)void register_task(const char *task_type,vector[int] params)TaskGraphResult generate_task_graph(int num_gpus, int my_gpu_id)vector[CppKNOperator*] operators3.4 C 代码muGraph在c代码中体现为mirage::kernel::Graph类这是最高层次的计算图。namespace mirage {namespace kernel {class Graph {private:struct pair_hash {size_t operator()(std::pairint, int const p) const;};public:Graph(dim3 gpu_dim {1, 1, 1}, bool disable_fingerprint false);~Graph();Graph(Graph const ) delete;Graph operator(Graph const ) delete;// input operatorDTensor new_input(std::vectorint const dims,std::vectorsize_t const strides,mirage::type::DataType data_type,mirage::layout::DmemLayout layout);DTensor elementunary(DTensor const input,mirage::type::KNOperatorType _type);// 忽略其它函数public:std::vectormirage::kernel::KNOperator * operators; // 操作符列表dim3 gpu_dim;off_t dmem_data_offset, dmem_fp_offset;std::vectorstd::pairoff_t, size_t allocated_data_tensors,allocated_fp_tensors;// Fields for persistent kernelsstd::mapmirage::type::GuidType, mirage::runtime::IODesc io_config;std::unordered_mapmirage::kernel::KNOperator const *,std::tupleint, int, runtime::TaskType, inttask_config;using OpType KNOperator;using TensorType DTensor;};mirage::kernel::Graph的主要特征是操作符类型使用KNOperatorType 枚举定义操作类型。张量表示使用DTensorDevice Tensor表示数据。操作节点包括输入KN_INPUT_OP输出KN_OUTPUT_OP矩阵乘法KN_MATMUL_OP等。mirage::kernel::Graph的成员函数以 elementunar 为例代码如下DTensor Graph::elementunary(DTensor const input,mirage::type::KNOperatorType type) {KNOperator *op create_elementunary_op(input, type);assert(op ! nullptr);operators.push_back(op);assert(op-output_tensors.size() 1);DTensor output op-output_tensors[0];return output;}3.5 KNOperatorGraph包含多个KNOperator对象。KNOperator是内核级别的操作符基类用于表示计算图中的节点。作为计算图中每个操作的基本单元可以维护输入和输出张量的信息提供操作类型表示。而且通过输入输出张量的连接关系可以建立操作间的依赖关系为后续的任务调度和事件管理提供基础。在runtime.cc中系统通过遍历Graph中的operators来生成任务图。class KNOperator {public:KNOperator(Graph *graph, mirage::type::KNOperatorType _type);KNOperator(Graph *graph,mirage::type::KNOperatorType _type,DTensor const input1);KNOperator(Graph *graph,mirage::type::KNOperatorType _type,DTensor const input1,DTensor const input2);KNOperator(Graph *graph,mirage::type::KNOperatorType _type,std::vectorDTensor const inputs);int get_input_dtensors(DTensor **inputs);int get_output_dtensors(DTensor **inputs);virtual ~KNOperator();virtual bool fingerprint(void) 0;virtual operator json() const 0; // 将操作序列转换为JSON格式// hash related functionsvirtual size_t get_owner_independent_hash() const;public:Graph *kgraph; // 通过该指针维护与所属计算图的关联mirage::type::KNOperatorType op_type; // 标识操作类型std::vectorDTensor input_tensors; // 存储操作的输入张量std::vectorDTensor output_tensors; // 存储操作的输出张量};KNCustomizedOpKNInputOpKNOutputOp是KNOperator的派生类。KNOperator的派生类举例。class KNInputOp : public KNOperator {public:KNInputOp(Graph *_graph,std::vectorint const dims,std::vectorsize_t const strides,mirage::type::DataType data_type,mirage::layout::DmemLayout layout,int3 input_map {-1, -1, -1});~KNInputOp();bool fingerprint(void);operator json() const override;public:std::vectorsize_t input_strides;int3 input_map;};class KNOutputOp : public KNOperator {public:KNOutputOp(Graph *_graph,DTensor const A,std::vectorsize_t const strides,int3 output_map {-1, -1, -1});~KNOutputOp();bool fingerprint(void);operator json() const override;public:std::vectorsize_t output_strides;int3 output_map;};class KNCustomizedOp : public mirage::kernel::KNOperator {public:KNCustomizedOp(Graph *_kgraph,std::vectorDTensor const inputs,mirage::threadblock::Graph const _graph);virtual ~KNCustomizedOp();bool fingerprint(void);size_t get_owner_independent_hash() const override;operator json() const override;public:mirage::threadblock::Graph bgraph;void get_bgraph(mirage::threadblock::Graph **bgraph);};KNOperatorType 的全量为enum KNOperatorType {KN_UNKOWN 1000,KN_INPUT_OP 1001,KN_OUTPUT_OP 1002,KN_MATMUL_OP 1003,// ElementUnaryKN_EXP_OP 1100,KN_SQUARE_OP 1101,KN_SQRT_OP 1102,KN_MUL_SCALAR_OP 1103,KN_SILU_OP 1104,KN_SIGMOID_OP 1105,KN_GELU_OP 1106,// non-lax elementunary opsKN_RELU_OP 1150,KN_CLAMP_OP 1151,KN_LOG_OP 1160,// ElementBinaryKN_ADD_OP 1200,KN_MUL_OP 1201,KN_DIV_OP 1202,KN_POW_OP 1203,// Reduction NormalizationKN_REDUCTION_0_OP 1300,KN_REDUCTION_1_OP 1301,KN_REDUCTION_2_OP 1302,KN_RMS_NORM_OP 1350,// Concat SplitKN_CONCAT_FIRST_OP_ID 1400,KN_CONCAT_0_OP 1400,KN_CONCAT_1_OP 1401,KN_CONCAT_2_OP 1402,KN_CONCAT_LAST_OP_ID 1409,KN_SPLIT_FIRST_OP_ID 1420,KN_SPLIT_0_OP 1420,KN_SPLIT_1_OP 1421,KN_SPLIT_2_OP 1422,KN_CHUNK_0_OP 1423,KN_CHUNK_1_OP 1424,KN_CHUNK_2_OP 1425,KN_SPLIT_LAST_OP_ID 1429,// CommunicationKN_ALLREDUCE_OP 1900,KN_CUSTOMIZED_OP 1999,};3.6 生成样例Kernel block图的生成逻辑如下从输入节点出发以xyz输入张量为起点初始化一个空前缀。迭代增长枚举算子来构造新节点每次枚举一个算子加入枚举matmul、add、exp...合成算子当枚举到合成算子马上进入block graph的synthesis每次扩张会检查合法性形状、显存/SMEM容量、路径约束。抽象剪枝计算当前前缀的抽象表达式E当和canonical form E0不一致时剪枝生成结束后会得到没有thread graph的kernel/block图候选集合。下面代码中给出了kernel graph和block graph的生成样例。import mirage as midef new_kernel_graph():kgraph core.CyKNGraph()return KNGraph(kgraph)def get_rms_linear():graph mi.new_kernel_graph() # kernel graphX graph.new_input(dims(num_tokens, 4096), dtypemi.float16)W graph.new_input(dims(4096, n_local_heads * head_dim 2 * n_local_kv_heads * head_dim), dtypemi.float16)# block graphtb_graph mi.new_threadblock_graph(grid_dim(384,1,1), block_dim(128,1,1), forloop_range32, reduction_dimx64)tX tb_graph.new_input(dtensorX, input_map(-1, -1, -1), forloop_dim1)tW tb_graph.new_input(dtensorW, input_map(1, -1, -1), forloop_dim0)tM tb_graph.matmul(tX, tW)tAccX tb_graph.forloop_accum(tX, rms)tAccM tb_graph.forloop_accum(tM)tO tb_graph.div(tAccM, tAccX)