英伟达CUDA太难！OpenAI出手要取代它

本文转载自量子位

用CUDA为GPU编程实在太难了。

为了让没有CUDA编程经验的人写出和专家效率相当的GPU代码，现在OpenAI推出了一种新的语言和编译器——Triton。

它的难度比CUDA低，但是性能却可与之相媲美。

OpenAI声称：

Triton只要25行代码，就能在FP16矩阵乘法shang上达到与cuBLAS相当的性能。

OpenAI的研究人员已经使用Triton，来生成比同等Torch效率高出1倍的内核。

Triton项目的负责人Philippe Tillet说：“我们的目标是使Triton成为深度学习CUDA的可行替代方案。”

25行代码实现最佳性能

Triton起源于Tillet在2019年学术会议MLPF上的一篇论文，当时他还是哈佛大学的一名研究生。

Tillet解决的问题是如何开发一种cuDNN更具表现力的语言，既能够处理神经网络中涉及的矩阵的各种操作，同时兼具可移植性且以及和cuDNN相媲美的性能。

现代GPU大致分为三个主要组件——DRAM、SRAM、ALU，对这些资源进行调度管理十分复杂，即便是熟悉CUDA的程序员。

Triton可以将这些优化过程完全自动化，让开发者可以更好地专注于并行代码的高级逻辑。

以矩阵乘法为例，能够为逐元素运算和归约编写融合内核很重要，但考虑到神经网络中矩阵乘法任务的重要性，这还不够。

Triton非常适合这些应用，只需约25行Python代码即可实现最佳性能。

@triton.jit
def matmul(A, B, C, M, N, K, stride_am, stride_ak, 
            stride_bk, stride_bn, stride_cm, stride_cn,
            **META):
    # extract metaparameters
    BLOCK_M, GROUP_M = META['BLOCK_M'], META['GROUP_M']
    BLOCK_N = META['BLOCK_N']
    BLOCK_K = META['BLOCK_K']
    # programs are grouped together to improve L2 hit rate
    _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 (resp. rn) denotes a range of indices
    # for rows (resp. col) of C
    rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    # rk denotes a range of indices for columns 
    # (resp. rows) of A (resp. B)
    rk = tl.arange(0, BLOCK_K)
    # the memory addresses of elements in the first block of
    # A and B can be computed using numpy-style broadcasting
    A = A + (rm[:, None] * stride_am + rk[None, :] * stride_ak)
    B = B + (rk [:, None] * stride_bk  + rn[None, :] * stride_bn)
    # initialize and iteratively update accumulator
    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 level matrix multiplication
        acc += tl.dot(a, b)
        # increment pointers so that the next blocks of A and B
        # are loaded during the next iteration
        A += BLOCK_K * stride_ak
        B += BLOCK_K * stride_bk
    # fuse leaky ReLU if desired
    # acc = tl.where(acc >= 0, acc, alpha * acc)
    # write back result
    C = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)
    mask = (rm[:, None] < M) & (rn[None, :] < N)
    tl.store(C, acc, mask=mask)

而另一方面，在CUDA中实现类似的过程需要花费更多的精力，甚至可能会降低性能。

手写矩阵乘法内核的一个重要优点是它们可以根据需要进行定制，以适应其输入和输出的融合变换。

如果没有Triton，对于没有特殊GPU编程经验的开发者来说，矩阵乘法内核的修改是非常困难的。

Triton背后的原理

Triton 的良好性能，来自于以Triton-IR为中心的模块化系统架构，这是一种基于LLVM的中间表示。

@triton.jit decorator通过遍历提供Python函数的抽象语法树（AST），产生的Triton-IR使用通用SSA构建算法上的动态。

生成的IR代码随后由编译器后端进行简化、优化和自动并行化，然后转换为高质量的LLVM-IR（最终转换为 PTX）。

研究人员发现，数据可以通过查看计算密集型块级操作（例如tl.dot）的操作数自动存储到共享内存中，并使用标准活性分析技术进行分配/同步。

另一方面，Triton程序可以通过同时执行不同的内核实例跨SM进行高效和自动并行化，以及通过分析每个块级操作的迭代空间，并在不同的SIMD中进行充分分区将SM内单元并行化。

目前Triton仅适用于英伟达GPU，但官方表示AMD GPU以及CPU的版本正在开发