TileLang 实战用 IKET 为 CUDA Kernel 插桩并在 Perfetto 中剖析【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang本篇指南围绕 TileLang 仓库中的 IKETexperimental instrumentation tool示例展开介绍如何为 TileLang 生成的 CUDA 内核注入命名标记marker、warp 级范围range与可选运行时标量载荷runtime payload并借助外部 IKET 剖析器采集 trace、在 Perfetto 时间线中可视化。读完本文你将能独立运行 examples/iket 下的三个示例、采集并服务一份完整的 trace同时理解 IKET 会话、缓存行为与源码级插桩原理。IKET 是什么为 TileLang CUDA 后端打造的实验性剖析通道IKET 是 TileLang 中一套实验性的插桩工具面向targetcuda后端。它向生成的 CUDA 源码注入三类事件即时标记marker、warp 局部范围range以及可选标量载荷payload由外部 IKET 剖析器收集并导出可供 Perfetto 检查的 trace。IKET 是CUDA 工具不属于 TileLang 语言命名空间使用时单独导入import tilelang.language as T from tilelang.tools.cuda import iket这一集成基于 TileLang 常规的targetcuda后端不依赖 TileScale 或 CuTe DSL 前端。仓库中 IKET 的 Python 实现位于 tilelang/tools/cuda/iket分为五个模块session.py编译会话与回调生命周期、frontend.pymark/range/payload等内核侧事件助手、codegen.pyCUDA 源码插桩、metadata.py随 TIR 传递的事件元数据与cli.py路径助手与剖析命令构造。环境要求与验证运行 IKET 示例需要满足支持 CUDA 的 TileLangCUDA 显卡与驱动运行示例所需的 PyTorch外部 IKET Python 包与运行时。用一条命令验证当前 Python 环境python -c import tilelang, iket, torch从代码生成路径看见 codegen.py 中_cluster_rank_instructionTileLang 同时支持 pre-Hopper 与 Hopper 目标SM90 及以上读取%cluster_ctarank寄存器记录集群 rankpre-Hopper 或无法确定架构时回退为 rank0不发射 Hopper 专用寄存器。三个示例文件与覆盖范围examples/iket 下有三个脚本覆盖 IKET 的完整功能面文件覆盖内容minimal.py即时标记iket.mark与一个 warp 局部范围iket.rangepayload_minimal.py单个int32运行时载荷标记all_features.py范围、无载荷标记、运行时载荷标记全特性三个脚本都以iket.session(...)包裹内核构造与编译并包含正确性自检如 minimal.py 通过torch.testing.assert_close(c, a b)验证结果保证插桩不会破坏计算语义。直接运行示例编译并运行最小标记示例python examples/iket/minimal.py \ --iket-output-dir /tmp/tilelang_iket_minimal运行带运行时载荷捕获的最小载荷示例python examples/iket/payload_minimal.py \ --iket-output-dir /tmp/tilelang_iket_payload_minimal \ --iket-runtime-payloads运行全特性示例python examples/iket/all_features.py \ --iket-output-dir /tmp/tilelang_iket_all_features \ --iket-runtime-payloads这三个命令会校验内核正确性并落盘生成的 CUDA 源码源码分别写为tilelang_iket_cuda_tool_kernel.cu、tilelang_iket_payload_minimal_kernel.cu、tilelang_iket_all_features_kernel.cu路径由--iket-output-dir决定。--iket-output-dir的默认值来自环境变量TL_IKET_OUTPUT_DIR未设置时各脚本有各自的/tmp/...默认目录。脚本还会在控制台打印event_table、runtime_payloads_enabled、instrumented检查生成的 CUDA 源码是否包含__iket_meta_info与TL_IKET_EVENT等诊断信息。注意直接运行只编译、执行插桩内核要收集 trace 需通过外部 IKET 剖析器运行程序见下节。采集一条 trace全特性示例对应的完整剖析命令rm -rf /tmp/tilelang_iket_all_features_profile python -m iket.cli.main \ --output-dir /tmp/tilelang_iket_all_features_profile \ --clobber \ profile \ --postprocess all \ -- \ python examples/iket/all_features.py \ --iket-output-dir /tmp/tilelang_iket_all_features_profile \ --iket-runtime-payloads剖析器在--之后启动目标命令并为其配置外部 IKET 运行时。成功后在输出目录中会出现四类文件iket_pid_0x....pftrace iket_pid_0x....pftrace.gz iket_pid_0x....trace.json iket_pid_0x....htmlTileLang 还提供了在 Python 中构造同一 shell 命令的助手command iket.profile_command( [python, examples/iket/all_features.py, --iket-runtime-payloads], directory/tmp/tilelang_iket_all_features_profile, ) print(command)profile_command(...)只返回经过引号处理的命令字符串不会自行启动剖析器。用 Perfetto 查看与检查 trace服务剖析输出目录使生成的 HTML 能加载其相邻的 trace 文件cd /tmp/tilelang_iket_all_features_profile python3 -m http.server 8080然后打开生成的精确文件名http://localhost:8080/iket_pid_0x....html在远程主机上可先做端口转发ssh -L 8080:localhost:8080 userremote-host如果页面只显示 Perfetto 着陆页则需在 Perfetto UI 中手动导入对应的.pftrace文件。下图是 all_features.py 产生的一条典型 trace蓝色主条带为block_total范围每个 block 一次warp 粒度记录其内叠加各 mark 事件与store_index等运行时载荷标记底部还有内核级事件。trace 的 JSON 导出可编程检查例如提取所有store_index标记的payloadValimport json from pathlib import Path trace_path max( Path(/tmp/tilelang_iket_all_features_profile).glob(*.trace.json), keylambda path: path.stat().st_size, ) data json.loads(trace_path.read_text()) launch data[launches][0] names data[stringTable] store_indices [ marker[payloadVal] for marker in launch[markers] if names[marker[markerNameIdx]] store_index and payloadVal in marker ] print(store_indices[:8])深入标记、范围与运行时载荷 API即时标记iket.mark在程序点发射一个即时事件iket.mark(load_inputs)范围iket.range与显式 push/pop用iket.range(...)作为 Python 上下文管理器包住词法区域with iket.range(compute): # TileLang 语句 ...显式范围 API 同样可用iket.range_push(compute) # TileLang 语句 iket.range_pop(compute)iket.range_start(...)/iket.range_end(...)分别是range_push/range_pop的别名见 frontend.py。范围起点可携带载荷范围结束事件不能携带载荷结束事件固定使用事件 ID31见 frontend.py。两个重要约束可从 frontend.py 的_get_event与_validate_payload_compat印证IKET按 warp 粒度记录范围一个包含四个 warp 的 block对同一个词法iket.range(...)会产生四条 trace 范围事件与范围名称限制为32 个 UTF-8 字节MAX_EVENT_NAME_BYTES 32见 metadata.py超长会抛出ValueError在同一前端注册表中同一名称的标记/范围不得混用不同载荷 dtype否则注册时报错。运行时载荷iket.payload标记与范围起点可捕获一个 32 位标量值TileLang 当前支持三种载荷 dtypeint32uint32float32对 TileLang 表达式使用显式载荷描述符最稳妥iket.mark(store_index, payloadiket.payload(i, dtypeint32)) iket.mark(scale, payloadiket.payload(value, dtypefloat32))简单 Python 标量bool/int/float或带dtype属性的表达式可直接传入由_infer_payload_dtype自动推断bool→uint32、int→int32、float→float32但显式 dtype 能消除 trace schema 歧义。三者的 IKET 内部类型 ID 映射为int32→5、uint32→6、float32→13见 frontend.py。运行时捕获是可选开启的with iket.session(runtime_payloadsTrue): program instrumented_add(1024) kernel tilelang.compile(program, targetcuda)不开启runtime_payloadsTrue时载荷 schema 仍编码在 TIR 元数据 token 中但生成的 IKET 元数据声明为NoPayload、事件不写载荷值——普通标记记录保持 4 字节。开启后事件改为两次独立的 32 位写时间戳/事件记录 载荷值见 codegen.py 的TL_IKET_EVENT_PAYLOAD_U32宏。载荷通过 IKET 的 warp 级 dump 机制观察其值通常代表该机制选中的 lane而非 warp 内每个线程。深入编译会话与缓存行为iket.session(...)的完整签名与默认值with iket.session( reset_eventsTrue, overrideTrue, disable_on_exitTrue, output_dirNone, runtime_payloadsNone, disable_cacheTrue, ): ...各参数对全局状态的控制与 session.py 的_Session实现一一对应reset_events清空后续构建内核的前端事件分配frontend.reset()会清空事件表、范围表并重置事件 ID 计数已嵌入PrimFunc的元数据不受影响override允许 IKET 在最外层会话激活期间替换已存在的tilelang_callback_cuda_postproc回调disable_on_exit默认在退出时恢复先前注册的回调作用域内使用应保持默认output_dir创建目录、设置环境变量TL_IKET_OUTPUT_DIR并在会话期间配置 TileLang 的 IKET 路径助手runtime_payloads临时选择是否发射载荷值None表示沿用先前设置disable_cache默认绕过 TileLang 的KernelCache会话退出时恢复先前缓存开关状态。禁用内核缓存很关键尽管事件名称与 schema 已属于 TIR 缓存身份的一部分但回调激活与运行时载荷模式是宿主侧编译状态。复用未带回调编译的二进制、或在不同载荷模式下编译的二进制会产生缺失或过期的插桩。只有调用方完全掌控这些条件时才应设置disable_cacheFalse。CUDA 回调采用引用计数嵌套 IKET 会话保持外层回调活跃退出最外层会话时恢复 IKET 之前的回调输出目录、载荷模式与缓存状态也在会话后恢复_restore_state见 session.py。高级场景可使用底层生命周期助手iket.enable() iket.is_enabled() iket.disable() iket.enable_runtime_payloads() iket.runtime_payloads_enabled() iket.disable_runtime_payloads()日常使用优先iket.session(...)——即使编译抛出异常它的__exit__也会完成清理避免状态泄漏。深入插桩如何在编译中存活每次前端事件调用都会在 TIR 中携带一个规范元数据 token前缀__tl_iket_v1_base64url 编码 JSON见 metadata.py其中包含事件名称、种类、范围身份与载荷 schema。这带来两个重要结果预构建的PrimFunc在会话进出、前端注册表重置后仍保留事件元数据结构相似但事件名称不同的内核其 IR 缓存身份不同。会话激活期间tilelang_callback_cuda_postproc从生成的 CUDA 中恢复这些 token_canonicalize_cuda_events见 codegen.py为整个模块分配统一事件 ID发射 IKET 元数据数组__iket_meta_info与各__iket_evt_decl_*、__iket_range_decl_*声明并定义 NativeDump 事件宏。事件名称不会从进程局部的event_table()注册表恢复。无载荷事件写一条 32 位时间戳/事件记录载荷事件用两次独立的 32 位共享内存写分别记录时间戳与载荷值且载荷写是volatile的防止 ptxas 将其合并为STS.64——这种记录形态不被当前外部 IKET patcher 接受见 codegen.py 中的注释。输出助手与调试工具CUDA 工具附带少量宿主侧助手定义于 cli.pyiket.set_output_dir(/tmp/tilelang_iket) iket.output_dir() iket.output_path(kernel.cu) iket.trace_files() iket.profile_command([...], directory/tmp/tilelang_iket)trace_files(...)返回.trace.json文件按体积从大到小排序这些助手只管理路径与命令构造本身不采集 traceiket.event_table()返回近期构建内核时注册的事件仅用于检查不是代码生成的依据显式重置前端事件分配可调用iket.reset()。限制说明从文档与源码可以确认的当前边界仅支持 TileLang CUDA 后端运行时载荷仅限int32、uint32、float32事件与范围名称限 32 个 UTF-8 字节不生成源码位置表trace 中的locIdx是 IKET 运行时位置索引不是 Python 或 TIR 行号IKET 按 warp 粒度记录事件回调、载荷模式、事件注册表与内核缓存开关均为进程全局状态并发编译工作流需自行协调访问该集成依赖外部 IKET 运行时的私有元数据与 NativeDump 约定应视为实验性功能。故障排查载荷 schema 出现但没有payloadVal在启用运行时载荷的会话下编译内核with iket.session(runtime_payloadsTrue): kernel tilelang.compile(program, targetcuda)同时确认标记使用了受支持的载荷描述符iket.payload(expr, dtype...)。剖析器在 patch 载荷内核时失败载荷插桩要求两次独立的 32 位写。用nvdisasm检查生成的二进制nvdisasm kernel.cubin | grep -E STS|PMTRIG期望形态为STS [addr], timestamp_with_event_id STS [addr0x4], payload_value PMTRIG event_id若时间戳/载荷对出现STS.64说明插桩序列已不符合 IKET NativeDump 的 patch 约定原因与对策见上文“两次独立 32 位写”一节。扩展阅读完整 API、会话语义与限制说明见 IKET 剖析指南插桩实现源码session.py、frontend.py、codegen.py、metadata.py可直接运行验证的示例minimal.py、payload_minimal.py、all_features.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 实战：用 IKET 为 CUDA Kernel 插桩并在 Perfetto 中剖析
TileLang 实战用 IKET 为 CUDA Kernel 插桩并在 Perfetto 中剖析【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang本篇指南围绕 TileLang 仓库中的 IKETexperimental instrumentation tool示例展开介绍如何为 TileLang 生成的 CUDA 内核注入命名标记marker、warp 级范围range与可选运行时标量载荷runtime payload并借助外部 IKET 剖析器采集 trace、在 Perfetto 时间线中可视化。读完本文你将能独立运行 examples/iket 下的三个示例、采集并服务一份完整的 trace同时理解 IKET 会话、缓存行为与源码级插桩原理。IKET 是什么为 TileLang CUDA 后端打造的实验性剖析通道IKET 是 TileLang 中一套实验性的插桩工具面向targetcuda后端。它向生成的 CUDA 源码注入三类事件即时标记marker、warp 局部范围range以及可选标量载荷payload由外部 IKET 剖析器收集并导出可供 Perfetto 检查的 trace。IKET 是CUDA 工具不属于 TileLang 语言命名空间使用时单独导入import tilelang.language as T from tilelang.tools.cuda import iket这一集成基于 TileLang 常规的targetcuda后端不依赖 TileScale 或 CuTe DSL 前端。仓库中 IKET 的 Python 实现位于 tilelang/tools/cuda/iket分为五个模块session.py编译会话与回调生命周期、frontend.pymark/range/payload等内核侧事件助手、codegen.pyCUDA 源码插桩、metadata.py随 TIR 传递的事件元数据与cli.py路径助手与剖析命令构造。环境要求与验证运行 IKET 示例需要满足支持 CUDA 的 TileLangCUDA 显卡与驱动运行示例所需的 PyTorch外部 IKET Python 包与运行时。用一条命令验证当前 Python 环境python -c import tilelang, iket, torch从代码生成路径看见 codegen.py 中_cluster_rank_instructionTileLang 同时支持 pre-Hopper 与 Hopper 目标SM90 及以上读取%cluster_ctarank寄存器记录集群 rankpre-Hopper 或无法确定架构时回退为 rank0不发射 Hopper 专用寄存器。三个示例文件与覆盖范围examples/iket 下有三个脚本覆盖 IKET 的完整功能面文件覆盖内容minimal.py即时标记iket.mark与一个 warp 局部范围iket.rangepayload_minimal.py单个int32运行时载荷标记all_features.py范围、无载荷标记、运行时载荷标记全特性三个脚本都以iket.session(...)包裹内核构造与编译并包含正确性自检如 minimal.py 通过torch.testing.assert_close(c, a b)验证结果保证插桩不会破坏计算语义。直接运行示例编译并运行最小标记示例python examples/iket/minimal.py \ --iket-output-dir /tmp/tilelang_iket_minimal运行带运行时载荷捕获的最小载荷示例python examples/iket/payload_minimal.py \ --iket-output-dir /tmp/tilelang_iket_payload_minimal \ --iket-runtime-payloads运行全特性示例python examples/iket/all_features.py \ --iket-output-dir /tmp/tilelang_iket_all_features \ --iket-runtime-payloads这三个命令会校验内核正确性并落盘生成的 CUDA 源码源码分别写为tilelang_iket_cuda_tool_kernel.cu、tilelang_iket_payload_minimal_kernel.cu、tilelang_iket_all_features_kernel.cu路径由--iket-output-dir决定。--iket-output-dir的默认值来自环境变量TL_IKET_OUTPUT_DIR未设置时各脚本有各自的/tmp/...默认目录。脚本还会在控制台打印event_table、runtime_payloads_enabled、instrumented检查生成的 CUDA 源码是否包含__iket_meta_info与TL_IKET_EVENT等诊断信息。注意直接运行只编译、执行插桩内核要收集 trace 需通过外部 IKET 剖析器运行程序见下节。采集一条 trace全特性示例对应的完整剖析命令rm -rf /tmp/tilelang_iket_all_features_profile python -m iket.cli.main \ --output-dir /tmp/tilelang_iket_all_features_profile \ --clobber \ profile \ --postprocess all \ -- \ python examples/iket/all_features.py \ --iket-output-dir /tmp/tilelang_iket_all_features_profile \ --iket-runtime-payloads剖析器在--之后启动目标命令并为其配置外部 IKET 运行时。成功后在输出目录中会出现四类文件iket_pid_0x....pftrace iket_pid_0x....pftrace.gz iket_pid_0x....trace.json iket_pid_0x....htmlTileLang 还提供了在 Python 中构造同一 shell 命令的助手command iket.profile_command( [python, examples/iket/all_features.py, --iket-runtime-payloads], directory/tmp/tilelang_iket_all_features_profile, ) print(command)profile_command(...)只返回经过引号处理的命令字符串不会自行启动剖析器。用 Perfetto 查看与检查 trace服务剖析输出目录使生成的 HTML 能加载其相邻的 trace 文件cd /tmp/tilelang_iket_all_features_profile python3 -m http.server 8080然后打开生成的精确文件名http://localhost:8080/iket_pid_0x....html在远程主机上可先做端口转发ssh -L 8080:localhost:8080 userremote-host如果页面只显示 Perfetto 着陆页则需在 Perfetto UI 中手动导入对应的.pftrace文件。下图是 all_features.py 产生的一条典型 trace蓝色主条带为block_total范围每个 block 一次warp 粒度记录其内叠加各 mark 事件与store_index等运行时载荷标记底部还有内核级事件。trace 的 JSON 导出可编程检查例如提取所有store_index标记的payloadValimport json from pathlib import Path trace_path max( Path(/tmp/tilelang_iket_all_features_profile).glob(*.trace.json), keylambda path: path.stat().st_size, ) data json.loads(trace_path.read_text()) launch data[launches][0] names data[stringTable] store_indices [ marker[payloadVal] for marker in launch[markers] if names[marker[markerNameIdx]] store_index and payloadVal in marker ] print(store_indices[:8])深入标记、范围与运行时载荷 API即时标记iket.mark在程序点发射一个即时事件iket.mark(load_inputs)范围iket.range与显式 push/pop用iket.range(...)作为 Python 上下文管理器包住词法区域with iket.range(compute): # TileLang 语句 ...显式范围 API 同样可用iket.range_push(compute) # TileLang 语句 iket.range_pop(compute)iket.range_start(...)/iket.range_end(...)分别是range_push/range_pop的别名见 frontend.py。范围起点可携带载荷范围结束事件不能携带载荷结束事件固定使用事件 ID31见 frontend.py。两个重要约束可从 frontend.py 的_get_event与_validate_payload_compat印证IKET按 warp 粒度记录范围一个包含四个 warp 的 block对同一个词法iket.range(...)会产生四条 trace 范围事件与范围名称限制为32 个 UTF-8 字节MAX_EVENT_NAME_BYTES 32见 metadata.py超长会抛出ValueError在同一前端注册表中同一名称的标记/范围不得混用不同载荷 dtype否则注册时报错。运行时载荷iket.payload标记与范围起点可捕获一个 32 位标量值TileLang 当前支持三种载荷 dtypeint32uint32float32对 TileLang 表达式使用显式载荷描述符最稳妥iket.mark(store_index, payloadiket.payload(i, dtypeint32)) iket.mark(scale, payloadiket.payload(value, dtypefloat32))简单 Python 标量bool/int/float或带dtype属性的表达式可直接传入由_infer_payload_dtype自动推断bool→uint32、int→int32、float→float32但显式 dtype 能消除 trace schema 歧义。三者的 IKET 内部类型 ID 映射为int32→5、uint32→6、float32→13见 frontend.py。运行时捕获是可选开启的with iket.session(runtime_payloadsTrue): program instrumented_add(1024) kernel tilelang.compile(program, targetcuda)不开启runtime_payloadsTrue时载荷 schema 仍编码在 TIR 元数据 token 中但生成的 IKET 元数据声明为NoPayload、事件不写载荷值——普通标记记录保持 4 字节。开启后事件改为两次独立的 32 位写时间戳/事件记录 载荷值见 codegen.py 的TL_IKET_EVENT_PAYLOAD_U32宏。载荷通过 IKET 的 warp 级 dump 机制观察其值通常代表该机制选中的 lane而非 warp 内每个线程。深入编译会话与缓存行为iket.session(...)的完整签名与默认值with iket.session( reset_eventsTrue, overrideTrue, disable_on_exitTrue, output_dirNone, runtime_payloadsNone, disable_cacheTrue, ): ...各参数对全局状态的控制与 session.py 的_Session实现一一对应reset_events清空后续构建内核的前端事件分配frontend.reset()会清空事件表、范围表并重置事件 ID 计数已嵌入PrimFunc的元数据不受影响override允许 IKET 在最外层会话激活期间替换已存在的tilelang_callback_cuda_postproc回调disable_on_exit默认在退出时恢复先前注册的回调作用域内使用应保持默认output_dir创建目录、设置环境变量TL_IKET_OUTPUT_DIR并在会话期间配置 TileLang 的 IKET 路径助手runtime_payloads临时选择是否发射载荷值None表示沿用先前设置disable_cache默认绕过 TileLang 的KernelCache会话退出时恢复先前缓存开关状态。禁用内核缓存很关键尽管事件名称与 schema 已属于 TIR 缓存身份的一部分但回调激活与运行时载荷模式是宿主侧编译状态。复用未带回调编译的二进制、或在不同载荷模式下编译的二进制会产生缺失或过期的插桩。只有调用方完全掌控这些条件时才应设置disable_cacheFalse。CUDA 回调采用引用计数嵌套 IKET 会话保持外层回调活跃退出最外层会话时恢复 IKET 之前的回调输出目录、载荷模式与缓存状态也在会话后恢复_restore_state见 session.py。高级场景可使用底层生命周期助手iket.enable() iket.is_enabled() iket.disable() iket.enable_runtime_payloads() iket.runtime_payloads_enabled() iket.disable_runtime_payloads()日常使用优先iket.session(...)——即使编译抛出异常它的__exit__也会完成清理避免状态泄漏。深入插桩如何在编译中存活每次前端事件调用都会在 TIR 中携带一个规范元数据 token前缀__tl_iket_v1_base64url 编码 JSON见 metadata.py其中包含事件名称、种类、范围身份与载荷 schema。这带来两个重要结果预构建的PrimFunc在会话进出、前端注册表重置后仍保留事件元数据结构相似但事件名称不同的内核其 IR 缓存身份不同。会话激活期间tilelang_callback_cuda_postproc从生成的 CUDA 中恢复这些 token_canonicalize_cuda_events见 codegen.py为整个模块分配统一事件 ID发射 IKET 元数据数组__iket_meta_info与各__iket_evt_decl_*、__iket_range_decl_*声明并定义 NativeDump 事件宏。事件名称不会从进程局部的event_table()注册表恢复。无载荷事件写一条 32 位时间戳/事件记录载荷事件用两次独立的 32 位共享内存写分别记录时间戳与载荷值且载荷写是volatile的防止 ptxas 将其合并为STS.64——这种记录形态不被当前外部 IKET patcher 接受见 codegen.py 中的注释。输出助手与调试工具CUDA 工具附带少量宿主侧助手定义于 cli.pyiket.set_output_dir(/tmp/tilelang_iket) iket.output_dir() iket.output_path(kernel.cu) iket.trace_files() iket.profile_command([...], directory/tmp/tilelang_iket)trace_files(...)返回.trace.json文件按体积从大到小排序这些助手只管理路径与命令构造本身不采集 traceiket.event_table()返回近期构建内核时注册的事件仅用于检查不是代码生成的依据显式重置前端事件分配可调用iket.reset()。限制说明从文档与源码可以确认的当前边界仅支持 TileLang CUDA 后端运行时载荷仅限int32、uint32、float32事件与范围名称限 32 个 UTF-8 字节不生成源码位置表trace 中的locIdx是 IKET 运行时位置索引不是 Python 或 TIR 行号IKET 按 warp 粒度记录事件回调、载荷模式、事件注册表与内核缓存开关均为进程全局状态并发编译工作流需自行协调访问该集成依赖外部 IKET 运行时的私有元数据与 NativeDump 约定应视为实验性功能。故障排查载荷 schema 出现但没有payloadVal在启用运行时载荷的会话下编译内核with iket.session(runtime_payloadsTrue): kernel tilelang.compile(program, targetcuda)同时确认标记使用了受支持的载荷描述符iket.payload(expr, dtype...)。剖析器在 patch 载荷内核时失败载荷插桩要求两次独立的 32 位写。用nvdisasm检查生成的二进制nvdisasm kernel.cubin | grep -E STS|PMTRIG期望形态为STS [addr], timestamp_with_event_id STS [addr0x4], payload_value PMTRIG event_id若时间戳/载荷对出现STS.64说明插桩序列已不符合 IKET NativeDump 的 patch 约定原因与对策见上文“两次独立 32 位写”一节。扩展阅读完整 API、会话语义与限制说明见 IKET 剖析指南插桩实现源码session.py、frontend.py、codegen.py、metadata.py可直接运行验证的示例minimal.py、payload_minimal.py、all_features.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），仅供参考