尧图网站建设 
TileLang Carver 框架实战基于 Tile 结构的调度提示推荐引擎【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang导读Carver 是 TileLang 内置的一个轻量级调度提示推荐框架它通过融合硬件架构信息、用户定义的 tile 结构与内置启发式规则自动生成并排序面向矩阵乘、逐元素变换、归约类算子的 tile 配置tiling strategy / blocking scheme / scheduling hint。本文将以 tilelang/carver/README.md 为核心结合仓库源码讲解 Carver 的模板 API、Hint 数据结构、后端架构抽象以及如何将推荐结果适配到 Triton 等其他编译器帮助你在 GPU、CPU 与加速器后端上快速获得可落地的分块调度方案。一、Carver 是什么为编译器生成调度提示的框架在 TileLang 这类面向高性能内核的领域专用语言中循环的划分方式tile 结构直接决定内核在具体硬件上的性能表现。手工枚举分块参数block 大小、warp 数量、归约步长等既繁琐又难以覆盖不同设备的约束条件。Carver 正是为此设计它统一生成面向多后端的 tile 候选集并在生成时纳入硬件约束如 CUDA 共享内存容量smem_cap、warp 大小、CPU 缓存结构、可用张量指令等最终输出一组带排序的调度提示。从实现上看Carver 的核心 API 定义在 tilelang/carver/init.py 中对外暴露了架构抽象CUDA、CDNA、RDNA以及CPU、Metal操作模板MatmulTemplate、GEMVTemplate、ElementwiseTemplate、GeneralReductionTemplate、FlashAttentionTemplate底层调度生成器roller含policy、bestfit、rasterization等子模块。典型用法只需三步创建架构对象 → 构造模板 → 调用recommend_hints(topk)获取推荐配置。二、快速上手GeneralReductionTemplate 与 SSR 结构2.1 基础用法对于通用的循环嵌套Carver 提供了GeneralReductionTemplate。它接受一个由SSpatial空间轴与RReduce归约轴组成的结构字符串以及对应的各维形状from tilelang import carver from tilelang.carver.arch import CUDA # 实例化 RTX 4090 的 CUDA 设备对象 arch CUDA(nvidia/geforce-rtx-4090) # 描述如下循环嵌套 # for i in Spatial(1024): # for j in Spatial(1024): # for k in Reduce(1024): # ... carve_template carver.GeneralReductionTemplate( structureSSR, shape[1024, 1024, 1024], dtypefloat16, ).with_arch(arch) # 生成前 20 个 tile 候选即调度提示 hints carve_template.recommend_hints(topk20) for hint in hints: print(hint)输出示例截断如下{ block: [1, 128], thread: [1, 128], rstep: [64], ... }, { block: [2, 64], thread: [2, 64], rstep: [64], ... }, ... { block: [1, 16], thread: [1, 16], rstep: [512], reduce_thread: [8], ... }2.2 SSR 结构如何映射为计算GeneralReductionTemplate的initialize_function会按结构字符串逐轴解析S轴进入输出空间形状R轴被构造为te.reduce_axis最终通过te.compute构建一个带归约的 TVMPrimFunc见 tilelang/carver/template/general_reduce.py。源码中的关键校验包括structure与shape必须同时提供且长度一致shape各维必须为正整数结构字符串只允许S/R大小写均可否则抛出ValueError。一个由 S 与 R 组成的 tile 结构可以模拟大量场景SS表示二维逐元素操作SSR则可以表示一次通用的矩阵乘法。因此该模板是快速验证调度思想的通用入口。2.3 with_arch 与自动推断架构with_arch(arch)由基类BaseTemplate实现见 tilelang/carver/template/base.py它把架构写入模板的_arch字段并返回自身以支持链式调用。值得注意的是BaseTemplate的_arch字段默认通过auto_infer_current_arch自动推断也就是说即使不显式调用with_arch模板也会尝试探测当前运行环境的设备架构。recommend_hints(topk)本质上是对get_hardware_aware_configs(self._arch, topk)的封装后者由各子类实现并调用统一的get_roller_hints_from_func进入 roller 调度生成管线。三、MatmulTemplate矩阵乘的专用模板针对C A * B这类标准矩阵乘法Carver 提供了更精细的MatmulTemplate可自动推断线程块、warp 划分以及是否启用 Tensor Core 等策略from tilelang import carver from tilelang.carver.arch import CUDA arch CUDA(nvidia/geforce-rtx-4090) carve_template carver.MatmulTemplate( M1024, N1024, K1024, in_dtypefloat16, accum_dtypefloat16, out_dtypefloat16, ).with_arch(arch) # 获取描述该矩阵乘的符号化函数 func carve_template.equivalent_function() print(Equivalent Function:\n, func) # 生成提示 hints carve_template.recommend_hints(topk20) for hint in hints: print(hint)输出示例{ block: [32, 64], warp: [16, 32], rstep: [128], use_tc: True, ... }, { block: [64, 32], warp: [32, 16], rstep: [128], use_tc: True, ... }, ... { block: [256, 32], warp: [128, 16], rstep: [32], use_tc: True, ... }3.1 模板参数说明根据 tilelang/carver/template/matmul.py 的源码MatmulTemplate的完整参数如下参数类型默认值含义Mint必填矩阵 A 与 C 的行数Nint必填矩阵 B 与 C 的列数Kint必填A 的列数 / B 的行数归约维trans_AboolFalse乘法前是否转置 Atrans_BboolTrue乘法前是否转置 Bin_dtypestrfloat16输入矩阵数据类型out_dtypestrfloat16输出矩阵数据类型accum_dtypestrfloat16累加中间结果数据类型with_biasboolFalse是否叠加偏置项initialize_function内部要求M、N、K均为正整数否则断言失败并根据trans_A/trans_B计算输入与权重的实际形状input_shape (M, K)A 转置时为(K, M)weight_shape (K, N)B 转置时为(N, K)。计算完成后若with_biasTrue会追加C[i, j] Bias[j]若out_dtype ! accum_dtype则会插入一次类型转换节点。3.2 equivalent_function拿到可调度的符号函数equivalent_function()返回模板构建的PrimFunc即后续可供 roller 策略分析、也可直接交给调度器使用的符号化计算函数。这一能力让模板不仅是提示生成器同时还能作为内核的原型定义便于在 TileLang / TVM 生态内继续做进一步的调度变换。四、Hint调度提示的数据结构无论是GeneralReductionTemplate还是MatmulTemplate返回的每个 hint 都是一个Hint对象见 tilelang/carver/roller/hint.py其核心字段如下字段含义block线程块block各空间轴的 tile 大小thread不使用 Tensor Core 时各轴的线程划分warp使用 Tensor CoreMFMA/MMA时各轴的 warp 划分rstep归约轴的步长每次加载多少 K 数据reduce_thread归约轴上额外分配的线程数use_tc是否启用 Tensor Corevectorize各张量加载时的向量化宽度如{A_reindex: 8, B_reindex: 8}pipeline_stage软件流水线级数默认 1split_k_factorSplit-K 因子用于 SM 浪费优化TileLang 专属rasterization_plan光栅化block 映射策略output_strides输出张量的 stride 信息Hint.to_dict()在输出时会做精简use_tc为真时输出warp否则输出thread只有reduce_thread的乘积大于 1、vectorize非空、pipeline_stage ! 1等条件下才会带上相应字段。这也解释了为什么不同 hint 打印出来的键并不完全一致。五、支持的架构与扩展方式5.1 开箱即用的后端Carver 目前为以下后端提供开箱即用支持CUDA如arch CUDA(nvidia/geforce-rtx-4090)CDNAAMD GPU 类后端CPU另有RDNA与Metal架构类位于 tilelang/carver/arch/ 目录下。新增一种架构只需实现TileDevice的一个子类或提供自定义 target描述清楚以下约束即可共享/本地内存容量smem_cap、max_smem_usagewarp或向量大小warp_size缓存大小l2_cache_size_bytes等可用的张量指令available_tensor_instructions。TileDevice基类见 tilelang/carver/arch/arch_base.py统一定义了这些字段并声明了必须实现的get_avaliable_tensorintrin_shapes。5.2 CUDA 后端内部结构以下是 CUDA 后端的示意性代码节选自 tilelang/carver/arch/cuda.pyclass CUDA(TileDevice): def __init__(self, target: Union[tvm.target.Target, str]): ... self.platform CUDA # 设备约束 self.smem_cap device.max_shared_memory_per_block self.compute_max_core device.multi_processor_count self.warp_size device.warp_size ... self.transaction_size [32, 128] # 字节 self.bandwidth [750, 12080] # MB/s近似值 self.available_tensor_instructions None def get_avaliable_tensorintrin_shapes(self): self.available_tensor_instructions ( TensorInstruction(mma, [16, 16]), TensorInstruction(wmma, [16, 16]), ) return [t.shape for t in self.available_tensor_instructions] def __repr__(self): return fCUDA({self.target})在实际实现中CUDA构造器还会通过正则解析 SM 架构字符串如sm_90、sm_90a均解析为计算能力 90并据此暴露一系列能力判定函数is_volta_archsm 70–79、is_ampere_archsm 80–88、is_ada_archsm 89、is_hopper_archsm 90、has_mma_supportsm 80。每个代际支持的张量核心精度矩阵也被硬编码在源码中例如Volta(float16, float32)、(float16, float16)Ampere在 Volta 基础上新增(bfloat16, float32)、(int8, int32)、(int4, int32)等Ada进一步加入(float8_e5m2, float32)、(float8_e4m3, float32)Hopper与 Ada 一致。is_tensorcore_supported_precision(in_dtype, accum_dtype, arch)正是依据这些矩阵判断某组输入/累加精度是否支持 Tensor Core——这解释了MatmulTemplate输出中use_tc字段是如何被决定的。5.3 CDNA 与 CPUCDNAtilelang/carver/arch/cdna.py通过tvm.runtime.rocm(0)获取设备信息并针对 gfx950CDNA4 / MI350做了 160 KB LDS 的特殊处理若驱动报告的默认值小于 163840 字节则覆盖之CPUtilelang/carver/arch/cpu.py实现较为轻量注释指出 LLVM 后端本身无需精细调优仅保持接口一致性。六、将 Hint 适配到其他编译器以 Triton 为例Carver 推荐结果的一大价值在于跨编译器适配。假设拿到如下 hint{ block: [32, 64], warp: [16, 32], rstep: [128], use_tc: True, vectorize: {A_reindex: 8, B_reindex: 8} }在Triton中可以这样解读block_m 32, block_n 64, block_k 128潜在 warp 划分warp_m 16, warp_n 32vectorize加载数据时使用向量宽度 8若use_tc为真在支持的情况下优先使用 Triton 的TensorOpsTensor Core。这样即可快速测试多组配置而无需手工猜测参数组合。同样的思路可以推广到 TVM、TileLang 或其他领域专用编译器——Carver 输出的 Hint 是一份与后端无关的调度蓝图各后端只需定义自己的映射规则。七、支持的模板一览Carver 通过模板抽象了常见的循环模式当前内置模板包括模板适用场景核心构造参数GeneralReductionTemplate通用Spatial-Spatial-ReduceSSR等结构structure、shape、dtypeFlashAttentionTemplate类 Attention 操作带 flash 内存访问模式batch_size、num_heads、head_dim、seq_length、seq_kv_length、is_causal、各 dtypeMatmulTemplate标准矩阵乘C A * BM、N、K、trans_A、trans_B、三种 dtype、with_biasGEMVTemplatey Ax或y xA类操作N、K、trans_B、三种 dtype、with_biasElementwiseTemplate逐元素 / 逐点变换shape、dtype补充说明依据对应源码ElementwiseTemplatetilelang/carver/template/elementwise.py内部以B A 1构造计算图用于探索纯逐元素内核的最优分块与向量化GEMVTemplatetilelang/carver/template/gemv.py固定M 1描述一次矩阵向量乘输出形状为(N,)FlashAttentionTemplatetilelang/carver/template/flashattention.py将 QK^T 与 SV 两次矩阵乘建模为两个PrimFuncNode并通过边Edge串联成计算图再经TensorCorePolicy从输出节点统一发射配置仓库中还提供了卷积模板 tilelang/carver/template/conv.py可用于探索卷积类的调度提示。如果你的算子有独特的循环结构或约束完全可以仿照以上模板自定义新的专用模板例如针对卷积、Flash Attention 的变体等——只需实现initialize_function与get_hardware_aware_configs两个抽象方法见基类 tilelang/carver/template/base.py。八、底层原理roller 策略管线recommend_hints的最终执行路径收敛到 tilelang/carver/utils.py 的get_roller_hints_from_func它首先将模板持有的PrimFunc交给DefaultPolicy启发式最小化内存流量、最大化并行度随后尝试通过get_tensorized_func_and_tags识别可张量化Tensor Core的子图一旦识别成功便改用TensorCorePolicy发射含use_tc、warp等字段的配置。DefaultPolicy.emit_config见 tilelang/carver/roller/policy/default.py会经过计算基础 tile → 分配归约步长rstep→ DFS 枚举共享内存 tile 候选 → 依据内存流量、共享内存占用、每 SM 的 block 数、wave 数等指标排序的完整流程最终截取topk个最优结果。每个Hint还附带traffic、smem_cost、block_per_SM、num_wave、grid_size等评估量存于TileDict为后续筛选提供量化依据。此外roller 还内置了rasterizationblock 到 SM 的映射如 panel 光栅化与bestfitbest-fit 分配机制进一步细化 hint 的完整度。九、路线图与 TileLang 的端到端集成Carver 的官方 TODO 清单中明确列出了一项计划适配 tile language为 TileLang 提供现成的调度调用或封装器wrapper打通端到端集成。当前 Carver 已经作为 tilelang/carver/ 模块随 TileLang 源码发布但其产出仍以通用 Hint 字典为主未来若补齐面向 TileLang 自身调度语法的直接映射例如把block/warp/rstep/split_k_factor/pipeline_stage等字段翻译成 TileLang 的T.parallel、T.Pipelined、软件流水线等调度原语即可实现一次模板描述、自动产出可编译内核的完整链路。对于开发者而言当前阶段可以先将recommend_hints的结果作为搜索空间初始点配合 TileLang 的自动调优autotuning工具进一步精化。十、总结Carver 以模板描述算子结构 硬件约束建模 策略化搜索的方式把调度提示的生成从经验试错变成了可复现的工程流程统一 APIwith_arch(...)recommend_hints(topk)即可跨后端获取排序后的分块方案硬件感知TileDevice子类携带共享内存、warp、缓存、带宽与张量指令信息保证提示贴合实际设备结构可表达SSR等结构字符串与MatmulTemplate等专用模板覆盖从逐元素到注意力的大多数常见算子跨编译器可移植Hint 是无后端依赖的中间表示可映射到 Triton、TVM 与未来的 TileLang 调度封装。想深入阅读源码推荐从 tilelang/carver/template/base.py、tilelang/carver/roller/hint.py 与 tilelang/carver/utils.py 三个文件入手它们构成了模板层、数据结构层与策略层的完整闭环。【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 TileLang Carver 框架实战：基于 Tile 结构的调度提示推荐引擎
TileLang Carver 框架实战基于 Tile 结构的调度提示推荐引擎【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang导读Carver 是 TileLang 内置的一个轻量级调度提示推荐框架它通过融合硬件架构信息、用户定义的 tile 结构与内置启发式规则自动生成并排序面向矩阵乘、逐元素变换、归约类算子的 tile 配置tiling strategy / blocking scheme / scheduling hint。本文将以 tilelang/carver/README.md 为核心结合仓库源码讲解 Carver 的模板 API、Hint 数据结构、后端架构抽象以及如何将推荐结果适配到 Triton 等其他编译器帮助你在 GPU、CPU 与加速器后端上快速获得可落地的分块调度方案。一、Carver 是什么为编译器生成调度提示的框架在 TileLang 这类面向高性能内核的领域专用语言中循环的划分方式tile 结构直接决定内核在具体硬件上的性能表现。手工枚举分块参数block 大小、warp 数量、归约步长等既繁琐又难以覆盖不同设备的约束条件。Carver 正是为此设计它统一生成面向多后端的 tile 候选集并在生成时纳入硬件约束如 CUDA 共享内存容量smem_cap、warp 大小、CPU 缓存结构、可用张量指令等最终输出一组带排序的调度提示。从实现上看Carver 的核心 API 定义在 tilelang/carver/init.py 中对外暴露了架构抽象CUDA、CDNA、RDNA以及CPU、Metal操作模板MatmulTemplate、GEMVTemplate、ElementwiseTemplate、GeneralReductionTemplate、FlashAttentionTemplate底层调度生成器roller含policy、bestfit、rasterization等子模块。典型用法只需三步创建架构对象 → 构造模板 → 调用recommend_hints(topk)获取推荐配置。二、快速上手GeneralReductionTemplate 与 SSR 结构2.1 基础用法对于通用的循环嵌套Carver 提供了GeneralReductionTemplate。它接受一个由SSpatial空间轴与RReduce归约轴组成的结构字符串以及对应的各维形状from tilelang import carver from tilelang.carver.arch import CUDA # 实例化 RTX 4090 的 CUDA 设备对象 arch CUDA(nvidia/geforce-rtx-4090) # 描述如下循环嵌套 # for i in Spatial(1024): # for j in Spatial(1024): # for k in Reduce(1024): # ... carve_template carver.GeneralReductionTemplate( structureSSR, shape[1024, 1024, 1024], dtypefloat16, ).with_arch(arch) # 生成前 20 个 tile 候选即调度提示 hints carve_template.recommend_hints(topk20) for hint in hints: print(hint)输出示例截断如下{ block: [1, 128], thread: [1, 128], rstep: [64], ... }, { block: [2, 64], thread: [2, 64], rstep: [64], ... }, ... { block: [1, 16], thread: [1, 16], rstep: [512], reduce_thread: [8], ... }2.2 SSR 结构如何映射为计算GeneralReductionTemplate的initialize_function会按结构字符串逐轴解析S轴进入输出空间形状R轴被构造为te.reduce_axis最终通过te.compute构建一个带归约的 TVMPrimFunc见 tilelang/carver/template/general_reduce.py。源码中的关键校验包括structure与shape必须同时提供且长度一致shape各维必须为正整数结构字符串只允许S/R大小写均可否则抛出ValueError。一个由 S 与 R 组成的 tile 结构可以模拟大量场景SS表示二维逐元素操作SSR则可以表示一次通用的矩阵乘法。因此该模板是快速验证调度思想的通用入口。2.3 with_arch 与自动推断架构with_arch(arch)由基类BaseTemplate实现见 tilelang/carver/template/base.py它把架构写入模板的_arch字段并返回自身以支持链式调用。值得注意的是BaseTemplate的_arch字段默认通过auto_infer_current_arch自动推断也就是说即使不显式调用with_arch模板也会尝试探测当前运行环境的设备架构。recommend_hints(topk)本质上是对get_hardware_aware_configs(self._arch, topk)的封装后者由各子类实现并调用统一的get_roller_hints_from_func进入 roller 调度生成管线。三、MatmulTemplate矩阵乘的专用模板针对C A * B这类标准矩阵乘法Carver 提供了更精细的MatmulTemplate可自动推断线程块、warp 划分以及是否启用 Tensor Core 等策略from tilelang import carver from tilelang.carver.arch import CUDA arch CUDA(nvidia/geforce-rtx-4090) carve_template carver.MatmulTemplate( M1024, N1024, K1024, in_dtypefloat16, accum_dtypefloat16, out_dtypefloat16, ).with_arch(arch) # 获取描述该矩阵乘的符号化函数 func carve_template.equivalent_function() print(Equivalent Function:\n, func) # 生成提示 hints carve_template.recommend_hints(topk20) for hint in hints: print(hint)输出示例{ block: [32, 64], warp: [16, 32], rstep: [128], use_tc: True, ... }, { block: [64, 32], warp: [32, 16], rstep: [128], use_tc: True, ... }, ... { block: [256, 32], warp: [128, 16], rstep: [32], use_tc: True, ... }3.1 模板参数说明根据 tilelang/carver/template/matmul.py 的源码MatmulTemplate的完整参数如下参数类型默认值含义Mint必填矩阵 A 与 C 的行数Nint必填矩阵 B 与 C 的列数Kint必填A 的列数 / B 的行数归约维trans_AboolFalse乘法前是否转置 Atrans_BboolTrue乘法前是否转置 Bin_dtypestrfloat16输入矩阵数据类型out_dtypestrfloat16输出矩阵数据类型accum_dtypestrfloat16累加中间结果数据类型with_biasboolFalse是否叠加偏置项initialize_function内部要求M、N、K均为正整数否则断言失败并根据trans_A/trans_B计算输入与权重的实际形状input_shape (M, K)A 转置时为(K, M)weight_shape (K, N)B 转置时为(N, K)。计算完成后若with_biasTrue会追加C[i, j] Bias[j]若out_dtype ! accum_dtype则会插入一次类型转换节点。3.2 equivalent_function拿到可调度的符号函数equivalent_function()返回模板构建的PrimFunc即后续可供 roller 策略分析、也可直接交给调度器使用的符号化计算函数。这一能力让模板不仅是提示生成器同时还能作为内核的原型定义便于在 TileLang / TVM 生态内继续做进一步的调度变换。四、Hint调度提示的数据结构无论是GeneralReductionTemplate还是MatmulTemplate返回的每个 hint 都是一个Hint对象见 tilelang/carver/roller/hint.py其核心字段如下字段含义block线程块block各空间轴的 tile 大小thread不使用 Tensor Core 时各轴的线程划分warp使用 Tensor CoreMFMA/MMA时各轴的 warp 划分rstep归约轴的步长每次加载多少 K 数据reduce_thread归约轴上额外分配的线程数use_tc是否启用 Tensor Corevectorize各张量加载时的向量化宽度如{A_reindex: 8, B_reindex: 8}pipeline_stage软件流水线级数默认 1split_k_factorSplit-K 因子用于 SM 浪费优化TileLang 专属rasterization_plan光栅化block 映射策略output_strides输出张量的 stride 信息Hint.to_dict()在输出时会做精简use_tc为真时输出warp否则输出thread只有reduce_thread的乘积大于 1、vectorize非空、pipeline_stage ! 1等条件下才会带上相应字段。这也解释了为什么不同 hint 打印出来的键并不完全一致。五、支持的架构与扩展方式5.1 开箱即用的后端Carver 目前为以下后端提供开箱即用支持CUDA如arch CUDA(nvidia/geforce-rtx-4090)CDNAAMD GPU 类后端CPU另有RDNA与Metal架构类位于 tilelang/carver/arch/ 目录下。新增一种架构只需实现TileDevice的一个子类或提供自定义 target描述清楚以下约束即可共享/本地内存容量smem_cap、max_smem_usagewarp或向量大小warp_size缓存大小l2_cache_size_bytes等可用的张量指令available_tensor_instructions。TileDevice基类见 tilelang/carver/arch/arch_base.py统一定义了这些字段并声明了必须实现的get_avaliable_tensorintrin_shapes。5.2 CUDA 后端内部结构以下是 CUDA 后端的示意性代码节选自 tilelang/carver/arch/cuda.pyclass CUDA(TileDevice): def __init__(self, target: Union[tvm.target.Target, str]): ... self.platform CUDA # 设备约束 self.smem_cap device.max_shared_memory_per_block self.compute_max_core device.multi_processor_count self.warp_size device.warp_size ... self.transaction_size [32, 128] # 字节 self.bandwidth [750, 12080] # MB/s近似值 self.available_tensor_instructions None def get_avaliable_tensorintrin_shapes(self): self.available_tensor_instructions ( TensorInstruction(mma, [16, 16]), TensorInstruction(wmma, [16, 16]), ) return [t.shape for t in self.available_tensor_instructions] def __repr__(self): return fCUDA({self.target})在实际实现中CUDA构造器还会通过正则解析 SM 架构字符串如sm_90、sm_90a均解析为计算能力 90并据此暴露一系列能力判定函数is_volta_archsm 70–79、is_ampere_archsm 80–88、is_ada_archsm 89、is_hopper_archsm 90、has_mma_supportsm 80。每个代际支持的张量核心精度矩阵也被硬编码在源码中例如Volta(float16, float32)、(float16, float16)Ampere在 Volta 基础上新增(bfloat16, float32)、(int8, int32)、(int4, int32)等Ada进一步加入(float8_e5m2, float32)、(float8_e4m3, float32)Hopper与 Ada 一致。is_tensorcore_supported_precision(in_dtype, accum_dtype, arch)正是依据这些矩阵判断某组输入/累加精度是否支持 Tensor Core——这解释了MatmulTemplate输出中use_tc字段是如何被决定的。5.3 CDNA 与 CPUCDNAtilelang/carver/arch/cdna.py通过tvm.runtime.rocm(0)获取设备信息并针对 gfx950CDNA4 / MI350做了 160 KB LDS 的特殊处理若驱动报告的默认值小于 163840 字节则覆盖之CPUtilelang/carver/arch/cpu.py实现较为轻量注释指出 LLVM 后端本身无需精细调优仅保持接口一致性。六、将 Hint 适配到其他编译器以 Triton 为例Carver 推荐结果的一大价值在于跨编译器适配。假设拿到如下 hint{ block: [32, 64], warp: [16, 32], rstep: [128], use_tc: True, vectorize: {A_reindex: 8, B_reindex: 8} }在Triton中可以这样解读block_m 32, block_n 64, block_k 128潜在 warp 划分warp_m 16, warp_n 32vectorize加载数据时使用向量宽度 8若use_tc为真在支持的情况下优先使用 Triton 的TensorOpsTensor Core。这样即可快速测试多组配置而无需手工猜测参数组合。同样的思路可以推广到 TVM、TileLang 或其他领域专用编译器——Carver 输出的 Hint 是一份与后端无关的调度蓝图各后端只需定义自己的映射规则。七、支持的模板一览Carver 通过模板抽象了常见的循环模式当前内置模板包括模板适用场景核心构造参数GeneralReductionTemplate通用Spatial-Spatial-ReduceSSR等结构structure、shape、dtypeFlashAttentionTemplate类 Attention 操作带 flash 内存访问模式batch_size、num_heads、head_dim、seq_length、seq_kv_length、is_causal、各 dtypeMatmulTemplate标准矩阵乘C A * BM、N、K、trans_A、trans_B、三种 dtype、with_biasGEMVTemplatey Ax或y xA类操作N、K、trans_B、三种 dtype、with_biasElementwiseTemplate逐元素 / 逐点变换shape、dtype补充说明依据对应源码ElementwiseTemplatetilelang/carver/template/elementwise.py内部以B A 1构造计算图用于探索纯逐元素内核的最优分块与向量化GEMVTemplatetilelang/carver/template/gemv.py固定M 1描述一次矩阵向量乘输出形状为(N,)FlashAttentionTemplatetilelang/carver/template/flashattention.py将 QK^T 与 SV 两次矩阵乘建模为两个PrimFuncNode并通过边Edge串联成计算图再经TensorCorePolicy从输出节点统一发射配置仓库中还提供了卷积模板 tilelang/carver/template/conv.py可用于探索卷积类的调度提示。如果你的算子有独特的循环结构或约束完全可以仿照以上模板自定义新的专用模板例如针对卷积、Flash Attention 的变体等——只需实现initialize_function与get_hardware_aware_configs两个抽象方法见基类 tilelang/carver/template/base.py。八、底层原理roller 策略管线recommend_hints的最终执行路径收敛到 tilelang/carver/utils.py 的get_roller_hints_from_func它首先将模板持有的PrimFunc交给DefaultPolicy启发式最小化内存流量、最大化并行度随后尝试通过get_tensorized_func_and_tags识别可张量化Tensor Core的子图一旦识别成功便改用TensorCorePolicy发射含use_tc、warp等字段的配置。DefaultPolicy.emit_config见 tilelang/carver/roller/policy/default.py会经过计算基础 tile → 分配归约步长rstep→ DFS 枚举共享内存 tile 候选 → 依据内存流量、共享内存占用、每 SM 的 block 数、wave 数等指标排序的完整流程最终截取topk个最优结果。每个Hint还附带traffic、smem_cost、block_per_SM、num_wave、grid_size等评估量存于TileDict为后续筛选提供量化依据。此外roller 还内置了rasterizationblock 到 SM 的映射如 panel 光栅化与bestfitbest-fit 分配机制进一步细化 hint 的完整度。九、路线图与 TileLang 的端到端集成Carver 的官方 TODO 清单中明确列出了一项计划适配 tile language为 TileLang 提供现成的调度调用或封装器wrapper打通端到端集成。当前 Carver 已经作为 tilelang/carver/ 模块随 TileLang 源码发布但其产出仍以通用 Hint 字典为主未来若补齐面向 TileLang 自身调度语法的直接映射例如把block/warp/rstep/split_k_factor/pipeline_stage等字段翻译成 TileLang 的T.parallel、T.Pipelined、软件流水线等调度原语即可实现一次模板描述、自动产出可编译内核的完整链路。对于开发者而言当前阶段可以先将recommend_hints的结果作为搜索空间初始点配合 TileLang 的自动调优autotuning工具进一步精化。十、总结Carver 以模板描述算子结构 硬件约束建模 策略化搜索的方式把调度提示的生成从经验试错变成了可复现的工程流程统一 APIwith_arch(...)recommend_hints(topk)即可跨后端获取排序后的分块方案硬件感知TileDevice子类携带共享内存、warp、缓存、带宽与张量指令信息保证提示贴合实际设备结构可表达SSR等结构字符串与MatmulTemplate等专用模板覆盖从逐元素到注意力的大多数常见算子跨编译器可移植Hint 是无后端依赖的中间表示可映射到 Triton、TVM 与未来的 TileLang 调度封装。想深入阅读源码推荐从 tilelang/carver/template/base.py、tilelang/carver/roller/hint.py 与 tilelang/carver/utils.py 三个文件入手它们构成了模板层、数据结构层与策略层的完整闭环。【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考