翻译 OpenAI Triton Programming Language for Neural Networks

翻译 OpenAI Triton Programming Language for Neural Networks

OpenAI Triton Programming Language for Neural Networks | Triton Programming Language | Analytics Steps
谁曾想过GPU编程会成为现实。但是，借助Nvidia的CUDA，它确实成为了现实。并行计算平台CUDA在2020年5月稳定发布后成为了热门话题。尽管它成为了每个开发者追求更好性能的首选，但其冗长的代码执行遭到了公开批评。

开发人员迫切需要一种开源编程语言，能够将繁琐的代码开发任务分解为简单、较小的片段。OpenAI在注意到CUDA的限制后，设计并发布了其开源编程语言Triton 1.0。

 什么是Triton？

Triton 1.0是一种类似Python的开源编程语言。这种新的编程语言可以重新定义编码过程。Triton专注于减少代码大小。它可以用更少的代码行执行与其他编程语言相同的功能。

它可以在25行以下编写F16矩阵乘法代码。使用Triton生成的内核比之前的torch实现更高效和更快。Triton可以为DNN转换器提供稳定的接口。

(Must Check: Working With Random Numbers in Python: Random Probability Distributions)

 Triton的优势

• Triton简化了特定内核的开发，比其他通用库更快。
• Triton的编译器简化代码，还自动化和并行化代码。编译器进一步将代码转换为可执行格式。
• 没有经验或者经验很少的开发者也可以在Triton中编写CUDA代码。
• Triton增加了以更少的努力达到硬件最高性能的可能性。 

具有较少GPU编程知识的程序员将通过Triton获得更好的前端平台。学习GPU编程的语法可能很容易，但是将算法移植以实现高效利用GPU并不容易。

在深入了解Triton的基础知识之前，让我们先了解一下GPU编程的困难之处。

GPU里面的DRAM，SRAM和ALU分别是什么，和CPU主机里的内存CPU什么关系？

在GPU（图形处理单元）中，DRAM（Dynamic Random Access Memory）、SRAM（Static Random Access Memory）和ALU（Arithmetic Logic Unit）是主要的硬件组件，而它们在功能和作用上与CPU主机里的内存和CPU紧密关联，但有着不同的特点。

1. DRAM（Dynamic Random Access Memory）：

   • 功能：用于存储程序和数据，是GPU的主存储器。
   • 特点：DRAM是一种动态随机存取内存，它是易失性的，需要不断刷新以保持数据的存储。
   • 关系：类似于CPU主机里的系统内存（RAM），用于存储正在运行的程序和相关数据。

2. SRAM（Static Random Access Memory）：

   • 功能：SRAM用于存储临时数据，速度比DRAM更快，但相对较昂贵。
   • 特点：相对于DRAM，SRAM是一种静态存取内存，不需要刷新操作，但它的面积和功耗较大。
   • 关系：在GPU中，SRAM通常用于缓存，类似于CPU主机中的高速缓存（L1、L2、L3缓存）。

3. ALU（Arithmetic Logic Unit）：

   • 功能：ALU执行算术和逻辑运算，是GPU中进行计算的核心部件。
   • 特点：ALU负责执行各种算术和逻辑操作，包括加法、减法、乘法、逻辑运算等。
   • 关系：类似于CPU主机中的ALU，用于执行计算指令。

与CPU主机里的内存和CPU的关系在于，GPU内部的DRAM类似于CPU主机中的系统内存，用于存储程序和数据。GPU的SRAM用于缓存，类似于CPU主机中的高速缓存。ALU负责执行计算，类似于CPU主机中的ALU。然而，GPU和CPU在体系结构上有很大的差异，GPU更专注于并行计算，具有大量的计算单元，而CPU更专注于单线程执行和更复杂的控制流程。因此，虽然有一些类似之处，但它们在设计和应用上存在明显的差异。

GPU编程的缺点

DRAM，SRAM和ALU是现代GPU编程的架构。在进行CUDA代码优化之前，必须考虑这三个因素。

• 内存从DRAM中传输必须集成到大型事务中，以利用现代内存接口的宽度。
• 数据在重新使用之前必须手动存储到SRAM中，为了最小化数据检索中的共享内存冲突，应高效地管理和利用数据。
• 将计算划分和调度到流多处理器中对于提升指令、线程级并行性非常重要。它还可以利用算术逻辑单元。

GPU也针对训练深度学习模型进行了优化。（深度学习是机器学习的一个子集，基于人工神经网络进行表示学习）。  

GPU可以同时处理多个计算，这一独特特性使其成为模型训练的基础。 GPU拥有极大数量的核心，使其能够为多个并行进程执行计算。

Triton专为GPU加速的神经网络而设计。神经网络和卷积神经网络是一系列算法，通过模仿人脑的处理过程来识别数据集之间的潜在关系。 在GPU上进行编程与深度学习相关。

Triton C编程语言的语法

Triton的编程结构基于ANSI C（CUDA C）。在CUDA C结构上进行了修改，以简化代码执行函数并使其能够适应不同的模型语义。

以下更改已集成到Triton中：

tile 声明：一种新的特殊语法将使多维数组的声明成为可能。示例（int tile[16],[16]）。特殊语法将突出嵌套数组的语义差异。在编写代码时，瓷砖形状必须保持不变。tile 形状可以通过tumble关键字进行参数化。可以使用省略号初始化一维整数瓷砖。

内置函数：（dot，trans，get_global_range）是一些内置函数，用于支持瓦片语义。Triton的基础设施还包括其他常用的智能数组操作（+，-，*，%）。

广播：可以使用（newaxis）关键字和slicing 语法来广播N维tile 。示例：int broadcast [9,9] = range [:,newaxis]

前缀“@”用于在tile 操作中执行控制流。

Triton的字符串和数字函数/操作

（请注意，字符串和数值函数的数据来自Triton）

1.triton.language.zeros
返回一个具有给定形状和数据类型的标量值为0的块。

• 形状：新数组的形状示例：（8，16），（8，0）
• 数据类型：新数组的数据类型

1. triton.language.broadcast_to
   尝试将给定的块广播到新的形状。

•  输入：输入块
•  形状：所需形状

1. triton.language.reshape 转换形状
   尝试将给定的块重新塑造成新的形状

•  输入：输入块
•  形状：所需形状

1. Triton.language.revel
   返回X的连续扁平化视图。

• X(Block)- 输入块
  5. triton.program_id
•  轴：value(0,1,2)
• builder:生成代码的IR构建器

1. Triton.language.num_programs

•  轴：值（0,1,2）
• 生成代码的IR构建器

1. triton.language.arange
   返回在开区间[start,end]内的连续值

• 开始：开始的区间幂次
• 停止：区间结束，二的幂次方
• 生成代码的IR构建器

1. triton.langauge.broadcast_to

•  输入：输入块
•  形状：期望的形状
• 生成器：用于代码生成的IR生成器
  9. triton.language.exp
• X(Block)：输入值
• 生成代码的IR构建器

 编程模型

CUDA GPU执行代码由SPMD编程模型支持。在SPMD编程模型中，每个内核（操作系统的核心部分）是启动网格中可识别的线程块的一部分。尽管Triton遵循类似的编程模型，但“triton”中的内核是“单线程的”（一次处理一个命令）。 

Triton中的内核自动并行化，并与不同的全局范围关联。自动并行化导致了更简单的内核，不存在类似CUDA的并发性。用于查询关联内核的全局范围的函数是get_global_range。

 Triton JIT 编译器

Triton IR使编译器能够自动执行基本优化。编译器中的数据会自动按操作数进行存储，以进行密集的块操作和同步分析技术。

Triton程序通过在SM之间并发执行不同的内核自动并行化，SM内的优化分析了SIMD单元（单一基础设施多数据处理单元）上的块级操作的迭代。

---

Triton中的自动并行化

Nvidia CUDA 对比 Triton

Triton和CUDA虽然有一些相似之处，但在许多方面也有不同之处。让我们了解一下这对的基本区别。

请注意，以下差异点来自OpenAI

现代GPU由三个主要组件组成：DRAM（动态随机存取存储器）、SRAM（静态随机存取存储器）和ALU（算术逻辑单元）。这三个组件都可以用于CUDA代码优化。

• 手动更新数据在SRAM中进行。在使用数据之前，SRAM管理数据重用并在数据检索过程中最小化共享内存冲突。
• DRAM内存传输与大型事务结合，以利用现代内存接口的优势。
• 计算在流多处理器（SMs）内进行排列和划分。多处理器为特殊目的ALU利用提供了线程级并行性的推动。

证明优化因素是困难的。即使是有着丰富CUDA经验的程序员也会在优化中遇到问题。Triton已经自动化了基本的性能级别优化，以减轻开发者的工作负担。开发者现在可以将注意力集中在并行代码的高级逻辑上。

CUDA Programming Model (Scalar Program, Blocked Threads)	Triton编程模型 （被阻止的程序，标量线程）
#pragma parallel for(int m = 0; i < M; m++) #pragma parallel for(int n = 0; j < N; n++){ float acc = 0; for(int k = 0; k < K;k ++) acc += A[i, k]* B[k, j]; C[i, j] = acc; }	#pragma parallel for(int m = 0; m < M; m += MB) #pragma parallel for(int n = 0; n < N; n += NB){ float acc[MB, NB] = 0; for(int k = 0; k < K; k += KB) acc += A[m:m+MB, k:k+KB] @ B[k:k+KB, n:n+NB]; C[m:m+MB, n:n+NB] = acc; }

 （来源：Triton）

	CUDA	OpenAI Triton
内存合并	手册	自动的
共享内存管理	手册	自动的
调度（在SM内部）	手册	自动的
调度（跨SM）	手册	手册

 （来源：OpenAI）

 Triton中的问题
Issues · openai/triton (github.com)

• 随机数生成
• 代码语法的小建议
• 点积计算错误的值
• 矩阵乘法示例中的PTX代码无效
• 非定义的常量表达式
• 性能差的matmul-square-nn测试