在AI系统领域，【【微信】】其实是个有些歧义的名字，因为至少有两个足够有影响力的【【微信】】相关的AI系统的工作，一个是NVIDIA用于在线服务布署的【【微信】】 Inference Server，另一个是由【【微信】】发起的高层次Kernel开发语言【【微信】】。

今天讨论的是后者。

大约在21年上半年，一个硬件公司的朋友给自己分享过【【微信】】的工作，当时并没有很深刻的映象。第二次关注是在2022年上半年，因为一些机缘巧合，又接触到了【【微信】】相关的内容，彼时【【微信】】基于MLIR重构的工作刚刚开始，甚至整个架子都还没能跑通。正好自己一直在思考基于MLIR解决计算密集算子codegen问题的可能性，也对【【微信】】这种高层次【【微信】】开发语言的想法有些好奇，于是对【【微信】】这个项目就有了更多实际的投入。随着参与的深入度，对【【微信】】也有了更多的一些理解。有些时候，在一些媒体上看到关于【【微信】】的一些说法，会觉得有些似是而非。缺失了技术细节的媒体讨论，很多时候，会给自己感觉是在想炮制一个吸引眼球的大新闻，而少了围绕事物客观度的本质探讨，反而带来不必要的负作用。严格来说，这篇文章也算是AI技术祛魅的一部分，不过【【微信】】并不仅限于大模型，所以就不再计入“大模型技术祛魅”的系列里。

【【微信】】在其社区网站上已经给出了【【微信】】的定位：

The aim of 【【微信】】 is to pro【【微信】】ironment to write fast code at higher producti【【微信】】, but also with higher flexibility than other existing DSLs.

我认为这句描述客观地反映了其定位。如果说【【微信】】在很多情况下要比基于CUDA层面写【【微信】】会更高效，是客观的（当然也存在一些稀奇古怪的【【微信】】通过【【微信】】开发非常别扭，而直接在CUDA层面开发更方便的case）。如果说【【微信】】写【【微信】】的效率又快，在性能上还总是能超过基于CUDA写的【【微信】】，就有些o【【微信】】了。实际上在上面这一句里提到的"higher flexibility than other existing DSLs"还是合适的，因为目前还没有看到和【【微信】】定位一样的DSL在开发【【微信】】这个方向能够做得比他更好。TVM算是一个类似的DSL，不过在描述【【微信】】的灵活性和效率上，【【微信】】相较TVM还是要胜出了不少，我会把这个原因归结为【【微信】】的定位小而美，不像TVM要解决那么宽阔的问题域，而是只解决计算密集型【【微信】】开发的问题，也没有一开始就花太大精力关注跨硬件的事情（虽然【【微信】】的设计架构允许对接多硬件，但其在多硬件上投入的资源相较TVM这样的工作还是少了不少）。另一个可能会被拿来比较的工作是CUTLASS。我的认识里，严格来说，CUTLASS不能算是一个DSL，而是一种提供了定制高性能【【微信】】开发的底层API，基于CUTLASS开发出来的【【微信】】的性能上界相较【【微信】】会高，但是在开发效率上，当前的CUTLASS相较【【微信】】还是有着明显的距离。

一言以概之，【【微信】】在【【微信】】的开发效率和性能之间选择了一个更适合上层用户的trade-off点。

下面会通过介绍【【微信】】的一些细节，来探讨【【微信】】是如何找到这个trade-off点的。

【【微信】】的工作，其实可以上溯到一个叫Isaac的工作，这个工作在2018年的GTC上曾经发过一篇poster，那一年我人在现场参加GTC，对这个工作还有一点模糊的映象。多年以后，现在的【【微信】】的工作和当年的isaac的工作之间还能看到联系，同样是根据workload的计算特点做problem-specific的tuning/codegen，只不过【【微信】】的工作在isaac的基础之上，加入了对【【微信】】硬件的支持，加入了对上层可编程DSL的支持，加入了【【微信】】的binding(这个对于AI生态很重要)，以及提供了相对丰富的预先实现好的优化【【微信】】集合。【【微信】】的作者Phil曾经在NV做过实习生，我想他从NV的实习经历里也应该获得了一些启发，再结合自己的经验，才在NV GPU的平台上创造出【【微信】】这个事物。一个工作，总需要保持一以贯之的投入和精进，并且不断吸收外界的反馈来作为演化的养分，才可能给这个领域带来更多的影响和帮助。【【微信】】的编程接口最开始其实是基于C++的API，这在【【微信】】 2019年的论文里还能看到痕迹。

对比一下实现相同功能，基于现在【【微信】】 【【微信】】 DSL的实现方式：

@triton.jit  def matmul(A, B, C, M, N, K,              **META):      BLOCK_M, GROUP_M = META['BLOCK_M'], META['GROUP_M']      BLOCK_N = META['BLOCK_N']      BLOCK_K = META['BLOCK_K']      _pid_m = tl.program_id(0)      _pid_n = tl.program_id(1)      pid_m = _pid_m // GROUP_M      pid_n = (_pid_n * GROUP_M) + (_pid_m % GROUP_M)       rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)      rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)       rk = tl.arange(0, BLOCK_K)       A = A + rm[:, None] + rk[None, :]      B = B +  rk [:, None] + rn[None, :]      # initialize and iterati【【微信】】r      acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)      for k in range(K, 0, -BLOCK_K):          a = tl.load(A)          b = tl.load(B)          # block le【【微信】】ion          acc += tl.dot(a, b)          A += BLOCK_K          B += BLOCK_K      C = C + (rm[:, None] + rn[None, :])      mask = (rm[:, None] < M) & (rn[None, :] < N)      tl.store(C, acc, mask=mask)  

从表述能力上来看，这段【【微信】】代码和上面帖出来【【微信】】原始论文里的C++代码是完全等效的。引入【【微信】】 DSL这件事情，从技术内核的角度对【【微信】】的帮助很有限，但是从【【微信】】被行业接受的角度来看却起到了很关键的作用。

本质上来说，【【微信】】和TVM/XLA这类工作的定位有所不同，如果说TVM/XLA是比较纯正的AI编译器的话，【【微信】】更像是一个面向AI加速器算子开发的领域开发语言，为了能够将用户使用【【微信】】语言开发的【【微信】】映射到具体硬件上的执行码，需要设计开发相应的【【微信】】 compiler来完成这层映射。所以当我们说【【微信】】的时候，其实隐指了【【微信】】语言+【【微信】】编译器这两个事物的综合体。

【【微信】】语言的语法并不复杂，在这篇文档里能够看到，大约只有50条左右的新增操作。当然【【微信】】是嵌入在【【微信】】这一宿主语言里，也复用了【【微信】】的一部分语言表达能力，比如上面例子里的for loop, 以及基本的 +=, , //等算术操作。【【微信】】从【【微信】】 DSL到【【微信】】底层IR的映射过程是通过解析【【微信】】 AST来完成的，所以并不是所有的【【微信】】写法都可以支持。并且不当的【【微信】】用法会带来类似于ONNX tracing那种所写非所得的歧义性。这其实是复用【【微信】】宿主语言一部分子集来简化新语言设计开发成本的DSL必然会遇到的问题了，比较经典的例子就是【【微信】】的设计对于对接新硬件以及脱离【【微信】】的布署侧带来的复杂性。不过到目前为止听到有用户抱怨【【微信】】这个问题的并不算多。我曾经思考过这个问题，感觉主要和几个原因有关:

• 【【微信】】的用户群和【【微信】】的用户群不完全一样，【【微信】】面向的用户属于工程能力相对较强的群体，并且【【微信】】的用户群也相较【【微信】】小众得多，因为【【微信】】毕竟处于比较底层的生态位。这类用户对于使用【【微信】】所能带来的收益（根据自己灵活的需求更快速地实现出性能还算体面的计算【【微信】】）和使用【【微信】】开发因为DSL语法限制带来的灵活性有着更强的包容度。
• 【【微信】】虽然是【【微信】】 DSL的使用界面，但其整体的使用流程更像TF 1.X，而不是【【微信】】的eager模式，所以如果在使用【【微信】】开发【【微信】】的时候出现问题，会在更早的阶段暴露出来，并反推用户进行改写适配，而不是代码写得满天飞，直到最后的某个集成环节才发现存在问题导致修改的feedback loop变长和成长增加。而【【微信】】的用户群的工程能力要求又使得他们对于开发过程中暴露出的问题相较普通AI建模人员有着更高的包容度。

【【微信】】的编程界面比较有意思的特点是，其操作的粒度大体来说既不是在标量这样细的颗粒度，也不是在Tensor这样相对高层次的颗粒度，而是在Block(或Tile)这种距离AI加速器更近的粒度(注: Block更像是一种逻辑概念，因为 其可以基于Tensor这一数据类型，配合一些属性完成定义)。一言以概，使用【【微信】】开发【【微信】】，Block这个层次的逻辑需要用户来处理，比如输入输出数据的分拆，Block粒度的数据累加等等。但Block内部的逻辑处理是交由【【微信】】编译器来自动完成的。具体到NV GPU硬件，SM(CTA)内部的逻辑细节用户不需要操心如何实现，会由【【微信】】来自动完成。这就相当大程度将用户从【【微信】】优化的细节中解放了出来，比如【【微信】】指令计算时将输入/输出数据分配到不同的计算线程的细节，以及为了减少shared memory访问通常会使用的swizzling的细节实现等。我们可以用一个相对简单的FP16的Matmul的例子来说明这一点。

【【微信】】的【【微信】】实现(在这个例子里为了简化描述：我把根据program_id对输入输出数据进行block拆分的逻辑给消除掉了，所以如果一个矩阵乘的计算不能在一个CTA里完成的话，严格来说这个【【微信】】是不完整的，不过这并不影响我们讨论【【微信】】的核心设计思想，所以后续不再解释)：

@triton.jit  def matmul_【【微信】】(      a_ptr, b_ptr, 

读创公司调研｜东港股份：已具备OpenAI的API接入能力