GPU 算子评测平台

写一个快的 GPU Kernel 很难，判断一个 Kernel 是不是真的快，同样很难。

普通 OJ 的判题逻辑很朴素：编译、喂输入、比对输出、看时间。但把判题对象换成 GPU 算子以后，几乎每一步都存在问题：

• 正确性不再是“完全相等”：浮点运算的结合顺序会影响结果，FP16 和 BF16 的误差量级差别很大，一个 Softmax 在大多数输入上都对，碰到一行全是 -inf 就输出了 NaN。
• 计时不再可信：GPU 调用是异步的，没有同步就停表，测到的只是下发时间；输入每次都一样，Kernel 可以偷偷缓存上一次的结果；数据刚好留在 L2 里，第二次测量天然更快；选手的代码和计时代码跑在同一个进程里，它甚至可以直接改掉计时函数。
• 一张卡不再只属于一个人：一块 GPU 很贵，大多数提交只需要几秒钟和几百 MB 显存，独占一张卡跑一次正确性检查过于浪费。但多个提交共享一张卡时，一个死循环的 Kernel 可能卡住整张卡，一次非法访存可能让同卡的其他进程一起报错，性能测试的数字也会因为别人在跑而剧烈抖动。
• 拉起不再是毫秒级：初始化 CUDA 上下文要几百毫秒，nvcc 编译要数秒，Triton 和 TileLang 第一次调用时还要 JIT，一个“只跑 3 毫秒”的 Kernel，从提交到出结果可能要等半分钟。

另一方面，今天给 GPU 写 Kernel 的已经不只是人了。越来越多的 Coding Agent 被用来自动生成和优化算子，而它们在“让评测器报出一个好看的数字”这件事上，往往比人更有创造力。已经有过公开的案例：某个自动写算子的 Agent 报告了数十倍的加速，事后发现是它利用了评测代码的漏洞，而不是真的写出了更快的 Kernel。

任务要求

本题希望你实现一个算子的在线评测平台（Kernel OJ）：用户提交一段算子实现，平台自动完成编译、多组用例的正确性检查和性能测量，给出可信的判题结果；同时，这个平台需要能够安全地运行不可信代码，让多个提交高效地共享有限的 GPU，并在高负载下保持可接受的响应速度。

提交的代码至少需要支持以下三种形式中的两种，推荐全部支持：

• CUDA C/C++：按照题目约定的函数签名实现，由平台编译并调用；
• Triton：提交 Python 代码，实现题目约定的入口函数；
• TileLang：同上。

如果你愿意，也可以额外支持 CuTe DSL、CUDA Python、Inline PTX 等其他形式，但需要说明它们为判题流程带来了哪些额外的问题。

技术选型不限，可适当部分使用开源项目。界面的美观程度不是本题考察重点，但平台需要具备 Lv1 中列出的基本功能，让老师和学生都能真正用起来。

单人完成本题时，完成 Lv1 和 Lv2 即可，Lv3-4 作为重要加分项；组队完成时需要完成 Lv1 到 Lv4，并在 Lv5 中至少选择一个方向。

Lv0：题库

在搭建平台之前，你需要先准备一个题库。请准备不少于 5 道题，每道题都应该是一个完整的题目包，至少包含：

• 题目描述、函数签名，以及三种语言各自的提交模板；
• 一个参考实现（通常是 PyTorch 的朴素实现），用作正确性的标准；
• 多组不同规模、不同数据分布的测试用例，其中应当包含边界情况，例如长度不是 Block 大小的整数倍、极大或极小的数值、全零、inf 和 nan。用例可以是预先生成好的输入与期望输出，也可以是带随机种子的输入生成器、在判题时与参考实现对比，或者两者结合，请说明你的选择；
• 正确性的判定方式：浮点结果按数据类型设置误差容忍度（atol / rtol），整数与排序类结果可能需要精确相等，随机类结果可能需要统计检验，请说明设置的依据；
• 性能测试使用的规模，以及一个性能基线（例如 PyTorch 原生算子或 cuBLAS 的耗时）；
• 一组样例提交：每道题至少有一份能通过的提交，以及一份错误的提交（如越界、竞争、精度不足）或一份正确但很慢的提交。所有题目的样例提交合计需要覆盖你所支持的每一种语言。

题库需要覆盖下面五类，每类至少一道：

注意，不同类别的题目会对判题提出不同的要求：浮点原子操作的累加顺序不固定，同一份代码两次运行的结果可能不同；排序和直方图的结果应当精确相等；蒙特卡洛类题目的输出本身带有随机性；FP64 的性能在数据中心卡和消费级显卡上相差可达数十倍。请在题目包中说明每道题的判题方式如何应对这些问题。

题目难度不必都很高，但至少有一道题应该让“正确”和“快”之间存在明显的距离。

题目包的格式请自行设计，但要能被平台直接加载。一个好的格式应该让出第 6 道题的人不需要读懂你的判题代码。

Lv1：判题闭环

Lv1 要求平台能够完成一次完整的判题：接收提交，编译，在多组用例上检查正确性，测量性能，并返回结果。

判题结果至少需要区分：通过、答案错误、编译错误、运行时错误（如非法访存）、超时、显存超限。当答案错误时，应当给出足够定位问题的信息，例如出错的用例规模、第一个超出容忍度的位置、期望值与实际值、最大误差，但不应泄露完整的隐藏用例。

正确性检查需要注意：

• 数值比较应当如何处理 NaN 和 Inf？参考实现本身输出了 NaN 时怎么办？
• 如何发现“大多数时候正确”的竞争条件？只跑一次够吗？

性能测量需要注意：

• 充分预热，并等待 GPU 上的异步任务完成后再计时；
• 多次重复，报告中位数而不是单次或最小值，并给出离散程度；
• 考虑 L2 缓存的影响，决定是否在每次测量前清空 L2，并说明理由；
• 考虑 GPU 频率的影响，例如 Boost 和降频带来的波动。

除了判题本身，平台还需要具备一个 OJ 的基本功能，至少包括教师端和学生端两种角色：

• 教师端：创建和发布题目，上传或编辑题目包，设置公开用例与隐藏用例、时间和显存限制、开放时间段；题目包在发布前应当经过自动校验，例如参考实现能正常运行、样例中的正确提交能通过、错误提交会被判错。老师还应当能够查看某道题的所有提交和整体通过情况。
• 学生端：浏览题目、下载提交模板、在线或通过命令行提交代码，查看每次提交的判题详情和自己的历史提交。
• 提交记录与来源：每一次提交都应当被完整记录，包括提交者、时间、语言、代码快照、判题结果、各用例的耗时，以及提交的来源（网页、命令行、API 或者 Agent）。任何一次判题都应当能被追溯和重新运行。
• 排行榜：按题目展示正确提交的性能排名。你需要考虑：排名依据哪个数字？两份提交的耗时只差 1%，而测量误差有 3%，谁排在前面？排行榜的成绩应当直接取自提交时的测量，还是在独占的卡上统一重测？排行榜展示的成绩是否会暴露隐藏用例的信息？

Lv1 结束时，请用 Lv0 中的样例提交跑通所有题目，并对同一份提交重复测量若干次，报告你的测量结果有多稳定。同时，请以老师的身份发布一道新题，再以学生的身份完成提交并出现在排行榜上，演示一次完整的流程。

Lv2：安全隔离

平台运行的是来自陌生人或者 Agent 的任意代码。Lv2 要求你先给出一份威胁模型，说明你需要防御什么、不防御什么，然后实现相应的隔离。

需要考虑的风险至少包括两类。

第一类是对系统的攻击。提交的代码可能读取宿主机上的文件、访问网络、Fork 炸弹、写满磁盘、在后台留下常驻进程，或者尝试利用 GPU 驱动和容器运行时的漏洞逃逸。CUDA 代码最终要以原生代码运行，Triton 和 TileLang 的提交本身就是任意 Python。你可以使用容器、gVisor、虚拟机、seccomp、cgroups 等手段，但需要说明：它们挡住了什么，没挡住什么，GPU 设备是如何被暴露给沙箱的，这条路径本身是否安全。

第二类是对判题的攻击。这是 Kernel OJ 特有的，也往往比第一类更隐蔽。例如：

• 读取参考实现或者预期输出；
• 直接调用 PyTorch 或 cuBLAS，而不是实现题目要求的算子；
• 根据输入的指针或内容缓存结果，第二次调用直接返回；
• 在另一个 Stream 上执行计算，让计时区间只覆盖到下发；
• 修改计时函数、同步函数或者比较函数；
• 只计算输出的一部分，赌判题只抽查了部分位置；
• 利用上一次运行残留在显存中的数据。

你需要让判题代码与选手代码的信任边界足够清晰。例如计时和比较是否应该与选手代码在同一个进程中？输入是否应该每次重新生成？哪些限制可以通过静态检查实现，哪些必须在运行时强制？

Lv2 要求你构造一套恶意提交的测试集，不少于 10 个，覆盖上面两类风险，并报告你的平台对每一个的处理结果。如果某些攻击仍然能够成功，请如实说明原因和可能的修补方式。

Lv3：共享一张卡

Lv3 要解决的问题是：让多个提交同时在一张 GPU 上运行，而且互不影响。

这里至少有三件事需要做到：

• 故障隔离：一个提交出现非法访存、死循环 Kernel、显存耗尽或者进程崩溃时，同卡的其他提交能够正常完成，平台能够发现并回收出问题的提交，GPU 在之后仍然可用。请分别测试上面几种故障，并说明哪些情况下仍然需要重置整张卡。
• 显存隔离：每个提交只能使用分配给它的显存额度，超出时应当得到明确的“显存超限”结果，而不是把别人挤出去或者让整张卡 OOM。
• 计算的计量与隔离：每个提交实际用了多少 GPU 时间、多少显存峰值、占了多少计算单元，需要被记录下来，用于计费、配额或者排查问题；同时，一个提交不应当能无限制地抢占其他提交的计算资源。

另外需要注意，正确性检查和性能测量对隔离的要求并不相同。正确性检查只关心结果对不对，适合高密度共享；而性能测量一旦与其他任务共享计算单元和显存带宽，结果就会失真。请量化这种干扰：同一份提交在独占和共享时，测得的耗时分别是多少，波动有多大。然后说明你的平台如何在资源利用率和测量可信度之间取舍。

Lv4：判题队列与拉起速度

当提交量上来以后，平台的瓶颈往往不在 GPU 上。Lv4 希望你把一次判题的耗时拆开来看，并把它压下去。

请首先将一次判题从提交到返回结果的耗时拆解成若干阶段，例如排队、沙箱启动、编译或 JIT、CUDA 上下文初始化、正确性检查、性能测量、结果回传，并测量每一段在冷启动和热启动下的耗时。然后针对最主要的开销进行优化。可以考虑的方向包括：

• 把编译放到 CPU 节点上，与 GPU 上的执行形成流水线；
• 预热一批已经完成初始化的沙箱或 CUDA 进程，提交到达时直接使用；
• 缓存编译产物和 JIT 结果。注意缓存不能在不同用户之间泄露代码或者被投毒；
• 用复用的沙箱时，如何确保上一个提交没有留下任何状态？

队列本身也需要设计：同一个用户在短时间内提交了一百次怎么办？比赛最后十分钟的提交洪峰怎么办？正确性检查和性能测量是否应该进入不同的队列？一份编译失败的提交是否还需要占用 GPU？

Lv4 需要给出压测结果：在一定的突发提交量下（例如 1 分钟内 200 次提交，混合不同题目和语言），报告判题延迟的 P50 和 P99、GPU 利用率，以及每一项优化分别带来了多少收益。

Lv5：让 AI 进场

Lv5 是可选方向，选择下面一个或多个即可：

• 面向 Agent 的判题接口：越来越多的选手是 Coding Agent。请为你的平台提供一套 Agent 友好的接口（例如 MCP 或者 CLI），并设计适合 Agent 阅读的反馈：结构化的编译错误、误差分布、Profiling 摘要（例如 Nsight Compute 给出的访存吞吐、占用率、瓶颈类型）。然后让至少一个 Coding Agent 在你的平台上从零开始完成两道题的多轮优化，对比“只给通过或不通过”与“给出详细反馈”两种条件下，Agent 最终达到的性能和消耗的轮数。注意，Profiling 工具通常需要读取硬件计数器，请说明你如何在沙箱中安全地提供它。
• AI 辅助出题与测例强化：给定一个 PyTorch 参考实现，能否让模型自动生成完整的题目包？更重要的是，能否自动生成“有杀伤力”的测例？一种检验测例强度的方式是变异测试：对一份正确的 Kernel 注入常见错误（越界、少一次同步、用错规约顺序、边界条件差一），看你的测例能杀死多少变异体。请报告 AI 生成测例前后的变异杀死率。
• AI 反作弊审查：Lv2 中的很多判题攻击难以用规则完全覆盖。请尝试让模型参与审查，例如判断一份提交是否真的实现了题目要求的算子，是否存在缓存结果、调用库函数、只计算部分输出等行为。你需要报告它的准确率和误报率：一份写得很奇怪但确实正确的高性能 Kernel，会不会被误判为作弊？它能否与运行时检测结合，而不是完全依赖模型的判断？

加分项

• Agent 代理访问：学生很可能让自己的 Coding Agent 代为读题、写代码和提交，也可能有完全自主的 Agent 直接来刷题。与 Lv5 中关注“反馈对 Agent 是否有用”不同，这里关注的是平台如何对待 Agent 这种访问者：
  • 身份与授权：Agent 代表哪个学生、能做什么？是否应该为 Agent 签发权限受限、可以随时吊销的凭证，而不是直接使用学生的账号密码？
  • 来源与规则：老师能否按题目或按作业设置“禁止 Agent”、“允许但需标注”、“完全开放”等规则？Agent 的提交在排行榜上是否应该与人类的提交分开？如果学生声称是自己写的，平台能否提供一些证据？
  • 配额与限流：Agent 可以不知疲倦地一小时提交上千次，把隐藏用例一点点试探出来，或者把判题队列挤满。如何为 Agent 设置配额，才能既不挡住正常的迭代，又不让它挤占其他同学的资源？
  • 机器可读：题目描述、提交模板、判题结果是否有稳定的结构化格式，让 Agent 不需要解析网页？题目描述中的内容会被 Agent 当作指令读取，老师或者其他用户能否借此对 Agent 进行提示注入？
• 测量的可信度：对同一份提交在不同时间、不同负载、不同 GPU 上的测量结果做一次系统的分析。除了绝对耗时，还可以给出每道题的理论下限（例如按访存量和峰值带宽估算的 Roofline），报告“距离硬件上限还有多远”，而不只是“比基线快多少”。
• 更细的错误诊断：集成 compute-sanitizer 等工具，在答案错误或运行时错误时，指出越界访问、未初始化读取或竞争条件的位置，并说明它带来的额外耗时是否值得。
• 多种计算卡：让平台同时管理不同类型的计算卡，包括但不限于不同代际的 NVIDIA GPU、昇腾 NPU、AMD GPU 等。支持哪些卡请根据你能拿到的硬件决定，但需要说明：
  • 题目包如何在不同硬件之间复用：同一道题的参考实现、误差容忍度和性能基线是否需要按硬件区分？Triton、TileLang 这类跨平台 DSL 的同一份提交能否直接在多种卡上判题？
  • Lv2 到 Lv4 的机制如何迁移：不同硬件的共享与隔离手段（如 NVIDIA 的 MPS / MIG、AMD 的计算分区、TPU 以芯片为单位的分配）、计时方式、故障恢复方式和拉起开销差别很大，你的平台在哪一层做了抽象，哪些部分只能针对单一硬件实现？
  • 同一份提交在不同硬件上的结果如何排名和展示？“在 A 卡上达到带宽上限的 80%”和“在 B 卡上比基线快 2 倍”能否放在一起比较？
• 可观测性：为平台提供一个面板，能看到每张卡上正在运行哪些提交、各自的显存和计算占用、队列长度，以及每一次判题各阶段的耗时。

思考题

• 用例应该像传统 OJ 一样预先存好输入和期望输出，还是每次随机生成输入、现场与参考实现对比？两种方式各自防得住哪些作弊，漏得掉哪些错误？一个 8192×8192 的 GEMM 用例要占多大的存储？参考实现在另一张卡上算出的期望输出还能直接用吗？随机出错的用例如何复现？
• 判题代码和选手代码应该在同一个进程里吗？分开以后，数据如何高效地在两者之间传递？
• 容器对于 GPU 来说算是足够的隔离吗？GPU 驱动是整个隔离链条中最薄弱的一环吗？
• 为什么性能测试要取中位数，而不是最小值？如果比赛以排名决定胜负，测量误差会带来什么问题？
• 如果选手的 Kernel 针对隐藏测例的规模专门调优，这算作弊吗？平台应该如何对待？
• 当 Agent 以每小时上千次的频率提交时，Kernel OJ 更像一个评测系统，还是一个强化学习的奖励函数？这对它的设计意味着什么？

这些问题没有标准答案。我们更关心你如何定义问题、做出取舍，并解释自己的设计。

附录：参考资料

背景阅读

• KernelBench：评估模型编写 GPU Kernel 能力的基准
• LeetGPU、Tensara：在线的 GPU 编程练习平台
• The AI CUDA Engineer：一个自动编写 CUDA Kernel 的 Agent，以及事后发现其评测被“钻空子”的讨论

Kernel 语言

• CUDA C++ Programming Guide
• Triton
• TileLang

其他计算卡

• ROCm / HIP，以及 AMD Instinct GPU 分区
• Pallas：JAX 上面向 TPU 和 GPU 的 Kernel 语言
• Triton-Ascend、TileLang Ascend

测量与诊断

• How to Implement Performance Metrics in CUDA C/C++
• triton.testing.do_bench：Triton 内置的计时方法，可以参考它对预热、L2 清空和统计的处理
• Nsight Compute
• Compute Sanitizer
• Roofline Model