65.9K
CodeProject 正在变化。 阅读更多。
Home

AMD GPU 的 GCN 汇编器

Ryan Scott White
starIconstarIconstarIconstarIcon
emptyStarIcon
starIcon

4.85/5 (17投票s)

2015 年 2 月 17 日

CPOL

46分钟阅读

viewsIcon

71296

downloadIcon

888

AMD 的 GCN (Generation Core Next Architecture) 汇编语言的汇编器/编译器

下载说明： 此项目需要 特定的硬件。（请参阅 要求 了解详情）

 

目录

简介
背景
Asm4GCN 汇编器子项目
   • 标签
   • 变量
         ◦ 声明 (已更新)
         ◦ 内联声明 
         ◦ 使用变量
         ◦ 重命名
         ◦ 变量索引 (已更新)
         ◦ 释放变量 (已更新)
         ◦ 强制寄存器编号 
    • 灵活的常量 
    • #define 支持
    • #s_pool / #v_pool 命令
    • 一行多条语句 & 分号 
    • 项目文件
    • 关注点
         ◦ 内联汇编 
         ◦ Smart 寄存器打包
OpenCLwithGCN 子项目
    • 特点 
         ◦ 内置文本模板引擎
         ◦ 混合 OpenCL 和 Asm4GCN 内核
    • 使用 OpenCLwithGCN
    • OpenCLwithGCN 项目文件
Asm4GcnGUI Windows 界面子项目
    • GUI 界面的特点
GPU 汇编通用话题
    • ​GPU 汇编的优缺点
    • 预加载的寄存器值
    • GCN 汇编代码编写技巧
限制
未来愿望清单
视频
系统要求
历史
其他 GCN 汇编器
感谢...

简介

这款 GCN GPU 汇编语言编译器可将人类可读的汇编转换为 AMD GCN 处理器上的二进制机器码。汇编语言是查看机器码的一种友好的抽象方式。除了一些例外，汇编语句大多直接映射到单个机器指令。对时间要求严格或需要使用特殊硬件功能的应用程序通常会使用汇编。GPU 汇编的一个应用实例可能是比特币挖矿应用程序。比特币挖矿是计算密集型的，需要微小的性能提升。使用汇编进行比特币挖矿还可以利用通常无法访问的特殊硬件指令或指令特性。此外，使用汇编编写代码通常可以创造性地组合指令。当然，使用汇编编写代码在可维护性和兼容性方面也有其缺点——详情请参阅 缺点 部分。

除了核心汇编器之外，此项目还包含另外两个 Visual Studio 项目。每个项目都以某种方式包装另一个项目。核心是 Asm4GCN 汇编器项目。在其外部是 OpenCLwithGCN 项目，它将汇编二进制文件注入到虚拟的 OpenCL 内核中。最后是 Asm4GCNGUI 项目，这是最外层。这个最终的 Windows 应用程序将所有内容整合到一个友好的语法高亮编辑器中。此编辑器还可以兼作 OpenCLwithGCN 的示例。

下面是每个项目的快速介绍，从汇编器核心开始，然后到用户界面。

Asm4GCN 汇编器 – 核心是 Asm4GCN 汇编器/编译器本身。它将汇编语句块转换为其二进制表示。它处理变量、预定义、标签和寄存器管理。它会向用户返回警告和错误。

OpenCLwithGCN – 这个项目接受一个包含 OpenCL 和 Asm4GCN 内核混合的 OpenCL 程序，然后注入（或修补）来自 gncAsm 内核的二进制文件。它还会运行程序。

Asm4GCNGUI – 这个简单的 Windows 应用程序实现了 Asm4GCN ，用于在一个小型编辑器中编写代码。该编辑器的语法高亮和代码补全功能使其便于编写 GCN 汇编。该编辑器可以快速测试 Asm4GCN 内核，并可能用于测试或教育目的。只需编写一些汇编和 OpenCL 内核以及基于 C# 的主机代码，然后观察结果。

背景

过去，我想利用编写汇编语言的额外性能优势，并利用只能通过汇编语言访问的特殊 GPU 硬件功能。为了满足 Nvidia GPU 的需求，我可以使用 Cuda 的 asm 函数内联一些 PTX。PTX 不是真正的汇编，但它提供了近乎真实体验。然而，对于 AMD，并没有真正直接的方法将汇编“内联”到 OpenCL 中。

最理想的解决方案是在 OpenCL 中提供一个 asm 函数，但目前这并不支持……至少在大多数情况下是不支持的。网上有一些帖子提到用户部分地实现了内联 asm，但有一个主要问题——无法传递变量。如果无法传递变量，则内联汇编用处不大。此外，使用的任何寄存器都可能覆盖一些正在使用的 OpenCL 变量。

由于我不知道有任何支持在 OpenCL 中内联 GCN 汇编的工具，所以我决定自己尝试。我最初的打算就是实现内联汇编，但效果不如我所愿——至少目前是这样。创建动态大小的虚拟内核并捕获使用的寄存器的复杂性很高，我的尝试并不可靠。经过大量努力，实现内联汇编的最初目标已被搁置，转而构建了一个允许 cl::program 包含 OpenCL __kernel 和 GCN __asm4GCN 内核的混合模式。实现内联汇编的框架已有一半可用，但我仍在为动态大小的虚拟内核捕获使用的寄存器方面苦苦挣扎。我不确定未来是否会添加内联汇编。（更多信息请参见 内联汇编 ）

我应该在这里补充一点，目前有一个名为 HetPas 的 Windows 应用程序，它生成可以在 OpenCL 程序中加载的 elf 图像文件。这个很棒的应用程序由 Realhet 创建，是一个闭源的 Windows 项目，它接受 GCN 汇编内核和类似 Pascal 的主机代码，然后运行它们——这与本项目类似。我使用 HetPas 已经有一段时间了，它工作得很好。我唯一觉得麻烦的是 Pascal 那部分。

Asm4GCN 汇编器子项目

此项目包含三个 Visual Studio 项目：Asm4GCN 汇编器、OpenCLwithGCN 和 Asm4GcnGUI。Asm4GCN 项目是此项目的核心汇编器，负责将一段汇编语句转换为原始二进制代码。

例如：[s_mov_b32 v1 s0] 被转换为硬件可读的 [7E020200]

除了将每个语句转换为其硬件代码对应项外，汇编器还处理标签和跳转语句等其他项目。我认为展示汇编器功能的最佳方式是通过流程图。为简洁起见，省略了一些步骤。

  1. 将一段 GCN 汇编代码传递给 Asm4GCN 汇编器进行二进制转换。
  2. 开始处理第一行（或下一行）。
  3. 删除此行上的任何空格。注释、空行和多余空格将被删除。
  4. 如果行以标签头（例如 myLabel:）开头，则在此处记录在标签列表中。
  5. #s/v_pool 和 #define
  6. 处理任何 #defines 标签。对于每个匹配项，关键字将被替换。
  7. 使用“;”将行分割成多条语句。
  8. 此行上还有（更多）语句吗？如果“是”，则解析该语句；如果“否”，则继续下一行。
  9. 该语句被分割成一个字符串数组。V_Mov v2,v4 被翻译成 {“V_Mov”, “v2”, “v4”}
  10. 处理任何变量初始化或销毁。在此，寄存器被保留或通过 free 返回到池中。
  11. 如果当前指令使用了任何变量，则用其当前的寄存器编号替换它们。
  12. 将十六进制、二进制、八进制和科学计数法字符串转换为常量。
  13. 如果当前指令引用了任何标签，则暂停转换指令为二进制，因为我们不知道常量值。
  14. 将当前指令转换为其二进制版本。
  15. 将当前指令添加到列表中以供后续处理。

 

下一步是为任何尚未填充 stmt.opSize 的语句填充它。完成此操作后，进行一些最终清理，然后创建一个二进制文件。

16. 创建一个具有未解析大小 (opSize) 的指令列表。这些是带有标签的指令，其距离未知。
17. 一条一条地处理每条指令，直到到达列表末尾。
18. 如果到标签的最小和最大距离相同，则指令大小必须相同。
19. 将字节大小设置为等同的最小/最大值，然后继续下一条指令。
20. 即使是跳跃距离最远的跳跃，字节大小仍然是 4 字节吗？如果是，则它将始终是 4 字节。
21. 设置 OpSize 为 4 字节
22. 我们已完成列表的枚举。是否有任何未解析的 OpSizes？如果是，我们将需要再次枚举。
23. 将任何带有标签的最终指令转换为二进制形式。
24. 将所有 4 字节和 8 字节指令合并成一个单一字节数组。
25. 返回完成的二进制文件。还会返回其他信息，例如寄存器使用情况。

标签

汇编器的主要任务之一是计算跳转标签的距离并将其转换为实际数字。正确地做到这一点是 GCN 汇编中更复杂的部分之一。事实证明，处理标签距离很像“先有鸡还是先有蛋”的困境。包含标签的指令可以是 32 位或 64 位，具体取决于跳转长度，但我们直到知道跳转和标签之间指令的大小才知道跳转长度。当跳转之间没有标签时，生活很简单，但一旦跳转越过其他跳转，我们就会遇到问题，因为我们不再知道确切的跳转距离。

为了解决这个问题，每条跳转语句都会尝试使用最短和最长的可能距离。最长距离是使用未解析指令的 64 位大小计算出的总大小。最短距离是相同的，但使用 32 位来表示未解析的指令。如果最短和最长距离都产生 32 位或 64 位指令，那么我们可以锁定该大小。我们一遍又一遍地对每个未解析的语句执行此操作，直到所有跳转语句都具有固定大小。一旦我们知道了每条指令的大小，就可以用适当的跳转长度常量替换标签。还有几个步骤，但这大致是思路。

作为上述方法的替代方案，一种简单但不那么高效的解决此问题的方法是，当不确定距离是否适合 32 位指令时，仅使用该指令的 64 位版本。64 位分支指令可以跳转 +/- 8TB 的距离！

标签系统的一些特点

  • 标签可以代替大多数常量。标签本身也是常量。它只是到另一条指令的距离，因此我们可以在通常使用常量的任何地方使用标签。例如，语句 s_mov s2, MyLabel 将简单地将 MyLabel 的距离移动到 s2。请记住，在这种情况下 s_sub s4, MyLabelA, MyLabelB 将失败，因为 s_sub 不能包含两个常量。
  • 支持具有相同名称的多个标签。对我自己来说，我经常想重用一个标签名称，或者用完我最喜欢的跳转标签。当使用具有相同名称的多个标签时，汇编器将使用最近的匹配标签。标签重用在复制粘贴代码块时可能很方便，因为通常不需要更新标签名称……在大多数情况下。不过，编译器会发出警告，因为标签名称重用有时会导致意外跳转。

变量

编写包含 10 到 20 条以上语句的汇编块会很快变得复杂。复杂性的一些原因在于试图记住每个寄存器做什么以及它是否可用。使用友好名称签出寄存器的能力非常有用。例如，名为 LoopCount 的变量比 S23:24 更具描述性。此外，签出寄存器的能力可以防止程序员意外覆盖现有寄存器。当可用寄存器数量变少时，花费大量时间来寻找可用寄存器可能会很耗时。变量将程序员从必须记住每个寄存器编号和可用寄存器中抽象出来。

示例

v4i myVar1, myVar2, mySum;
v_mov_b32 myVar1, 10;
v_mov_b32 myVar2, 20;
v_add_i32 mySum, myVar1, myVar2
free myVar1, myVar2

Asm4GCN 的变量系统包含三个部分——声明、使用变量和释放变量。

声明变量

声明变量时，请使用变量类型声明关键字，如 v4i、v4f、v8f、s8f 等。当编译器遇到声明时，它会立即从允许的寄存器池中保留第一个可用的寄存器。我试图使用一个简短的三位数声明来简化一个长的声明，如 scalar unsigned int 或 vector long long。第一个数字决定向量或标量。第二个数字是以字节为单位的大小。最后一个数字是变量的预期数据类型。类型后面是变量名本身。此字符必须是字母或下划线，其余字符可以包含数字。

声明示例

 v4i _my_int_vector;
 v8u myUnsignedLongVector;
 s8f myDoubleScalar1, myDoubleScalar2; //multiple items can be declared at once
 s4u myAddrForcedToS10 s10;     // forced to 10 (used for incoming params)
 v8f myDoubleForcedTo2 v[2:3];  // either v[2:3] or just the beginning, v2, can be specified

声明以类型开头，后跟一个或多个变量名。我想让变量类型保持简短和简单，所以我设计了一个三字符变量类型，包含类型、大小和数据类型。

第一个数字 

第一个数字指定变量将驻留在哪个内存空间。这是 S 表示标量或 V 表示向量内存。这会告诉寄存器保留系统保留一个标量或向量寄存器。

第二个数字 (已更新！)

中间数字指定变量的大小（以字节为单位）。大小告诉编译器保留一个或多个连续寄存器来保存变量，以及如何在内存中对它们进行对齐。大小值为 2 表示 2 字节，但会占用一个完整的 DWORD 大小的寄存器。值为 4 也会使用一个寄存器。目前，编译器对于 1、2 或 4 的处理没有区别——这仅供参考。但是，如果使用 8，这确实会影响编译器，因为它将保留两个连续的寄存器。多寄存器变量也已对齐——双寄存器变量将包含一个偶数寄存器后跟一个奇数寄存器。如果使用 16，将占用 4 个寄存器并按 4 对齐。任何大于 16 的值始终按 4 对齐。

按大小对齐和使用的寄存器
  使用的寄存器 对齐
1 字节 1 个寄存器 1 - 任何寄存器
2 字节 1 个寄存器 1 - 任何寄存器
4 字节 1 个寄存器 1 - 任何寄存器
8 字节 2 个寄存器 2 - 必须以 2 整除的寄存器开始
16 字节 4 个寄存器 4 - 必须以 4 整除的寄存器开始
32 字节 8 个寄存器 4 - 必须以 4 整除的寄存器开始
最后一个数字

最后一个数字是预期的数据类型。它是 F 表示浮点数，I 表示整数，U 表示无符号整数，或 B 表示位。 “位”格式实际上是一个占位符格式。位实际上可以表示任何其他格式，包括布尔值。不幸的是，最后一个格式说明符尚未被编译器使用，但它仍然为程序员提供了有用的信息。

我最初的打算是通过能够执行类似 R2 = R3 + R4 这样的简单命令来简化 GCN 汇编，其中需要类型信息。类型也可以用来发出警告。如果指定了浮点类型但与 v_add_i32 一起使用，则可能会发出警告。这些额外的功能将来可能会添加，也可能不会。

强制寄存器编号

声明寄存器时，可以选择将其强制指定为特定的寄存器。它可以分配一个物理寄存器编号，也可以重用过去变量的寄存器。

硬编码寄存器编号的两种方法

  1. 通过使用固定的寄存器编号。（例如 v4u myLaneID v0 ）
  2. 或者，通过使用过去的变量名。（例如 v4u myNewVar somePastVariable ）

固定寄存器编号通常用于在内核启动时捕获预先分配的寄存器。这些可以是参数地址、块 ID/大小或车道 ID 等项。具有固定寄存器的变量应在内核开始时使用。如果它在代码的后期使用，可能会失败，因为一个变量可能已经自动分配了该寄存器。

另一种分配寄存器的方法是复制过去变量的寄存器编号。这在某些情况下很有用。

有时寄存器的含义会发生变化，变量名不再准确。处理此问题的一种方法是“重命名”现有变量，但这可能很难看。这之前已实现但后来已删除。更好的方法是释放一个变量，然后在一个声明中使用相同的寄存器。这比笨拙的重命名语句能产生更清晰的代码。它也提供了更大的灵活性。

创建新变量时，程序员可能只想使用现有寄存器的一部分。例如，假设有一个变量占用了 4 个寄存器，但由于某种原因只需要保留最后一个寄存器。可以这样做：s4u myNewVar myPast16SizedVar[3]；现在 myPast16SizedVar 的前三个寄存器是免费的，最后一个寄存器可以保留一个不同的名称。可能会出现这样的问题：为什么不使用 s4u myNewVar; s_mov_b32 s4u myNewVar myPast16SizedVar[3]。有两个问题：(1) 会有 5 个寄存器的峰值使用量，(2) s_mov_b32 是硬件指令，因此会产生更大（更慢）的代码。之前的寄存器重用方法不占用任何指令——它只是更好的寄存器管理。

内联变量声明

编程时，尤其是在汇编中，寄存器声明会占用大量空间。我一直在思考如何解决这个问题，并回忆起 C 风格的实现方式。在 C 中，可以通过 int x = y + z; 来完成，但在调整了各种元素后，我得出了 v_add_i32 int x, y, z 这样的格式。

让我们通过一个实际示例来演示

首先，不带内联声明...

v4u vLocalSize
v_mov_b32     vLocalSize, localSize
v_mul_i32_i24 vLocalSize, groupId, vLocalSize 
v4u localSizeIdx
v_add_i32     localSizeIdx, vcc, laneId, vLocalSize 
v4u vGlovalID
v_add_i32     vGlobalID, vcc, baseGlobalId, localSizeIdx
v4u vGlobalOffset
v_lshlrev_b32 vGlobalOffset, 2, vGlobalID

现在带内联声明...

v_mov_b32     v4u vLocalSize, localSize
v_mul_i32_i24     vLocalSize, groupId, vLocalSize
v_add_i32     v4u localSizeIdx, vcc, laneId, vLocalSize 
v_add_i32     v4u vGlobalID, vcc, baseGlobalId, localSizeIdx
v_lshlrev_b32 v4u vGlobalOffset, 2, vGlobalID

哪个更清晰？我喜欢后者。

但等等，还有更多。由于寄存器可以在同一条指令上声明并自动释放，因此可以使用同一个寄存器。这可以通过减少寄存器数量来获得性能更好的代码。例如，在上面的代码中，v_add_i32 v4u vGlobalID, vcc, baseGlobalId, localSizeIdx 可以释放 localSizeIdx 并在同一行上重用同一个寄存器。它可能被翻译成类似 v_add_i32 v2, vcc, s8, v2. 的内容。请注意，v2 在同一条指令中被重用。

使用变量

只需在代码中使用变量名，而不是使用寄存器编号（例如 v17 或 s[20:21]）。在后台，每当汇编器遇到汇编中的变量时，它都会进行字典查找以检索其寄存器编号和类型。然后，它会将变量名文本替换为寄存器类型和编号。例如，v_add_i32 v3, v4, myInt; 变为 v_add_i32 v3, v4, v7;

变量索引

变量索引增加了访问多寄存器变量中特定寄存器的能力。例如，对于 s 8b myVar，可能需要访问第二个寄存器。这可以通过在变量后面附加 [1] 来实现。

在添加两个 64 位值时，可能需要这种方式

v8i myInt1, myInt2; // to be added together
[ myInt1 and myInt2 are assigned values here ]
v_add_i32 myInt1[0], vcc, myInt1[0], myInt2[0]  // note: "[0]" is actually redundant on this line
v_add_i32 myInt1[1], vcc, myInt1[1], myInt2[1]

释放变量

尽可能坚持自动释放。它能产生更清晰的代码，并提高寄存器利用率。

变量的寄存器可以通过自动释放或使用 free 关键字来释放。在后台，这会将变量标记为已完成，并将寄存器编号标记为空闲。在大多数情况下，自动释放就足够了。

自动释放

寄存器会在最后一次使用的语句上自动释放。编译器会记录变量使用的所有位置，并在最后一次使用的语句上自动释放该寄存器。自动释放还可以回收同一条语句中的寄存器，因此效率更高。

手动释放

使用 free 关键字后跟一个变量名来手动释放变量。

v4u myVar1, myVar2
...
v_add_u32 myVar1, myVar1, myVar2 <-- myVar1 and myVar2 last used (myVar2 auto freed here) 
..
free myVar1 <-- myVar1 is freed here

free 后面可以跟一个或多个用逗号分隔的变量名，例如 free myVar1, myVar2

这将在语句位置强制释放分配给它们的寄存器。Free 用于使变量保留的时间比自动释放时更长。在某些情况下这可能很有用。

一个原因是，当使用 GPR 索引通过 v_movrels 或 v_movreld 访问寄存器时。由于使用 v_movrel 指令不直接使用寄存器编号，自动释放不会知道它稍后在代码中使用，并会过早释放它。

第二个原因是让变量保留更长时间，因为有时我们会跳回代码中并期望一个变量仍然是活动的。如果一个变量在第 20 行最后一次使用后被自动删除，并且寄存器被第 21 行的新变量回收，然后我们稍后跳回到第 18 行，那么这些寄存器将被擦除。如果我们稍后在代码中添加一个 free 语句，那么编译器将使该变量及其寄存器保留更长时间。

灵活的常量

内联常量，也称为字面量，可与支持它们的指令一起使用。Asm4GCN 支持多种常量类型：十进制 (250)、十六进制 (0xFA)、八进制 (0x 372)、二进制 (0b11111010)、指数 (25E1, 25E+1) 和标签 (myLabel)。您可能在想为什么列出了标签？嗯，在 AsmGCN 中，标签只是到该标签的字节距离，因此它可以在任何常量位置使用。

有效示例

s_add_i32    s3, s4, 12             // pos int
s_add_i32    s3, s4, -12            // Small neg.
s_min_u32    s5, s6, 0xabcd         // Hex
s_min_u32    s5, s6, 10e2           // Exp
s_min_u32    s5, s6, -10e2          // Neg Exp
s_mov_b32    s4, 2.                 // float 
s_mov_b32    s4, -20.0              // neg float 
s_mov_b32    s4, .5                 // float 
s_mov_b32    s4, -.5                // float 
s_mov_b32    s4, 343.432            // float 
s_mov_b32    s4, 3.4e4              // exp float 
s_mov_b32    s4, -34.4e-4           // exp float 
s_mov_b32    s4, 0o7654             // Octal  (4000/FA0)
s_mov_b32    s4, 0b0011111111       // Binary (255/FF)

#define 支持

该汇编器支持 c 风格的 #define 语句。这通过在代码块中的每个 #define 进行简单的查找和替换来实现。例如，#define _fe_ 54321 会将代码中的所有 _fe_ 替换为 54321。使用下划线的原因是为了防止意外使用。如果我们只使用 fe ，我们可能会意外地将任何 fe 替换为 12345 ，从而使 v_bfe_u32 变成 v_b12345_u32。不带下划线的 #defines 仍然可以使用，但它们会生成编译器警告。

带参数的 #defines 也受支持，因此类似 #define _world_(AAA) 1AAA1 的内容将搜索 int myNum = _world_(00); 之类的项，并将它们转换为 int myNum = 1001;。

有效示例

#define _hw3_ Hello World!
#define _hw1_(opt0) Hello opt0 World!
#define _hw2_(opt0,opt1) Hello opt0 World from opt1!

#s_pool 和 #v_pool

大多数情况下，您可能希望使用寄存器 0 到 255，但有时可能需要使用特定集合（或池）的寄存器进行编译。#s_pool 和 #v_pool 可以放在汇编器文本的顶部来实现这一点。当遇到此编译器命令时，寄存器保留系统将使用这些寄存器作为变量。通常，这会放在代码顶部，但您实际上可以在任何地方使用它——如果它不在顶部附近，将出现警告。

有效示例

#S_POOL s22, s23, s24, s27, s29, s30, s31, s33, s34, s35, s36, s37
#V_POOL v11, v12, v13, v14, v15, v17, v19, v20, v21, v23, v24, v25

一行多条语句和分号 

在大多数情况下，一行可以使用多条语句。如果是这种情况，那么 分号 (;) 可以用来 分隔 语句。除 #v_pool、#s_pool 和 #define 外，所有指令和大多数命令都支持一行多条语句。对于只有单个命令或指令的行，可以 选择性地 使用分号。

示例

s8u myScalar; v8u myVector; v_add_i32 myVector, myVector, myScalar // this is okay
v_add_i32 v0, v1, v2  // this is okay
v_add_i32 v0, v1, v2; // this is okay
v_add_i32 v0, v1, v2; #define _myDef_ 12345678  // fail - #define must be on its own line

Asm4GCN 项目文件

以下是源代码文件及其描述的列表。行数显示在文件名之后（括号中）。该项目大约有 10,000 行（包括注释和空格）。

  • GcnISA.cs (1756) - 此文件包含有关 GCN 的所有原始数据。它包含指令信息、寄存器别名信息等的字典、数组和枚举。
  • Encoder.cs (1404) - 此文件包含静态 GcnParcer 类。该类包含每个 GCN 编码格式的方法，其任务是将单个语句行转换为其 OpCode 二进制格式。这是 GCN 汇编器的核心。
  • GcnBlock.cs (679) - GcnBlock.cs 包含 GcnBlock 类。该类负责将指令块转换为 byte[] 二进制形式。
  • DataStructs.cs (50) - 包含各种杂项结构，如 GcnStmt、Define 和 AsmVar。
  • Labels.cs (109) - 保存 Label 和 Labels 类。该类负责跟踪标签和跳转距离。
  • ParseOperand.cs (346) - 这个静态工具类负责解析操作数。它将十六进制、八进制和二进制字符串转换为常量，并验证数据类型是否允许。
  • Program.cs (208) - 包含 Asm4GCN 在命令行中使用时使用的功能。然而，大多数项目会像 DLL 一样链接函数。
  • RegPool.cs (356) - 使用 RegPool 类跟踪寄存器使用情况。RegPool 维护一个当前可用寄存器的数组。此类通过提供的允许寄存器列表或从零开始的范围进行初始化。
  • RegUsage.cs (111) - RegUsageCalc 类跟踪每个寄存器大小的使用计数。它还记录了它们出现的行号的最大值。这用于信息用途或了解内联语句可能需要多少寄存器。（内联尚未实现）
  • Log.cs (122) - 包含用于处理输出的日志记录类。它可以输出到 StringBuilder 或直接输出到控制台。
  • Tools.cs (76) - 包含一个静态扩展类，其中包含一些有用的扩展，例如 IsBetween()。
  • TestInput.txt (166) - 此文件包含 GCN 汇编的示例。它也用作测试。
  • Variables.cs (409) - 包含大多数与变量相关的类和任务。它也是 RegPool 和 RegUsage 实例的所有者。

关注点

内联汇编 – 内联汇编在这个项目中并没有真正奏效，但我想包含一些笔记，以防有人想尝试。我最初让它变得更复杂（此处未显示）并感到厌倦。“内联”代码仍然存在，但需要完成——由他人或我自己完成。这是一个计划的大纲：

  1. 定位并提取内联 Asm4GCN 块（到字符串中）
  2. 接下来，将提取的 Asm4GCN 块汇编为二进制，并记录 byteSize、sReg 和 vReg 计数。这是通过使用“OutputUsageMode”选项进行汇编来完成的。在此模式下，GCN 汇编被编译为临时寄存器，并记录每个大小和类型的寄存器（S 或 V）的数量。最大使用点很重要，因为。
  3. 在 OpenCL 内核中，我们需要用一些生成的虚拟 OpenCL 代码替换内联汇编。虚拟代码的字节大小需要与内联汇编二进制文件相同或略大。虚拟代码还需要使用相同数量的正确大小的 sReg 和 vRegs 以及相同的参数。它还必须由 barrier 和 mem_fence 包围，以防止 OpenCL 重新排序代码，否则虚拟代码以后将难以识别和替换。
  4. 现在将 OpenCL 虚拟代码编译为二进制，并记下使用的 sRegs 和 vRegs。了解使用的寄存器很重要，因为它告诉我们可以在汇编中使用哪些寄存器。编写内联汇编时，我们需要使用变量而不是固定寄存器编号，因为我们不知道将来会使用哪些寄存器池。
  5. 接下来，我们重新汇编内联汇编，但这次我们将在顶部添加一个允许的寄存器池，命令为 #S_POOL 和 #V_POOL。例如，#S_POOL s22,s23,s24,s27 将指示汇编器使用这些寄存器作为变量。
  6. 最后一步是定位整个程序二进制文件中的虚拟二进制文件，然后用完成的二进制文件替换它。可以使用 RegEx 来完成此操作。

未完成的部分是什么？步骤 (3) – 这有点奏效，但效果不如我预期的好。它确实创建了一个大小接近正确的内核，并且寄存器计数也接近，但并非完全准确。我的尝试的代码在 FillerKernelAttempts.cl: DummyFillerCode() 中。步骤 (6) – 这没有完成。动态创建的 DummyFillerCode 二进制文件的开头似乎很难找到。

智能寄存器打包/保留 – 为 GCN4Asm 开发的寄存器保留系统不像仅仅查找第一个可用寄存器那么简单。原因是根据选择的可用寄存器，可以容纳不同数量的寄存器。这就像将箱子装进卡车。如果箱子随意放置，会占用很多空间。但是，如果仔细选择箱子的位置，则可以装下更多箱子。

矩形打包用于二维，箱子打包，更受欢迎，是三维打包。

我进行了一些关于如何最好地分配寄存器的在线研究，但没有找到直接回答我问题的算法。我意识到寄存器保留非常类似于“矩形打包”。矩形打包是将小矩形装入大矩形的过程。使用此方法，小矩形可以代表变量，大矩形可以代表整个可用寄存器空间。A 寄存器的宽度将是小矩形之一的宽度，而寄存器的生命周期是矩形的高度。 下面是一个示例，其中每个字母代表一个变量的生命周期。

可视化寄存器使用示例
  R0 R1 R2 R3 R4 R5 R6 R7 R8 R9
指令 1 A A                
指令 2 A A B   C C C C    
指令 3 A A B   C C C C D D
指令 4     B   C C C C D D
指令 5     B           D D
指令 6 E   B   F F        
指令 7 E       F F        
指令 8 E   G G F F        

 

 

 

 

 

 

 

因此，使用类似于矩形打包的方法，我们只需要将这些小矩形（寄存器）装入一个大矩形中。同样，大矩形代表所有寄存器（宽度）以及代码块的整个生命周期（高度）。目标是装入所有小矩形并最小化大矩形的宽度。大矩形的宽度是我们总共使用的寄存器数量，因此我们希望尽可能减小它。使用的寄存器越少，我们将获得的占用率就越高。

然而，寄存器打包和矩形打包之间存在一些差异

  1. 在大多数情况下，寄存器生命周期的开始和结束是基本固定的且不可调整的，因此矩形不能上下移动，而在矩形打包中它们可以。但汇编中有一个例外，但这通常不是汇编器处理的任务。汇编语句通常可以重新排序或上下移动以最大程度地减少寄存器使用——参见 此处。
  2. 寄存器的宽度通常是二的幂（1、2、4、8、16...），而矩形打包通常允许任何宽度。也许这里有额外的优化机会。
  3. 最后，宽度为二的寄存器需要按二对齐（第一个寄存器必须是偶数），而宽度为四或更大的寄存器需要按四对齐。有一些额外的代码可以确保这一点。

尽管矩形打包并非完美契合，但它是计算机科学中一个知名且有据可查的领域，并且有一些著名的方法可以找到接近最优的解决方案。可能是最常见和最简单的方法是贪婪的大到小算法。我记得在我的计算机科学课上学过这个！该算法首先插入最大的项目，然后处理到最小的项目。它非常简单，但有效。

对于寄存器打包，我遵循了一个类似的系统，但不是按大小对块进行排序，而是将它们装入最小的空间。如果一个矩形完美地适合一个空槽，它就使用它，否则就将其装入它所能容纳的最小空槽中。该算法使用评分系统来决定这一点。

简而言之，GCN4Asm 比仅仅选择第一个可用寄存器空间更智能，但同时它也不是一个高度优化的算法。

OpenCLwithGCN 子项目

此项目使得可以在 C# 项目中使用 Asm4GCN。其主要功能是用虚拟 OpenCL 内核替换 Asm4GCN 内核，然后在汇编之后，用 Asm4GCN 内核二进制文件替换虚拟 OpenCL 二进制文件。OpenClwithGCN 项目允许在同一个 cl::program 中混合使用 Asm4GCN 和 OpenCL 内核。我认为用流程图来描述是最好的

OpenCLwithGCN 的特点

内置文本模板引擎

在汇编中编程几乎必须拥有某种机制来动态重复文本。汇编程序员的一个普遍需求是展开循环。在高级语言中，编译器会自动执行此操作，但在汇编级别，这由程序员负责。

此项目使用一个基于 C# 的文本模板转换引擎，可用于操作汇编。它由 C# 代码标签组成，在编译之前在汇编文本上运行——类似于预定义。对于 GCN 文本模板机制，我选择将 C# 代码包装在 [[..]] 样式的标签中。要使用它，只需在双括号内放入任何 C# 代码来控制显示内容的流程。变量可以使用 [[=myVar]] 打印在任何地方。更多功能可以在 https://codeproject.org.cn/Articles/867045/Csharp-Based-Template-Transformation-Engine 找到。

文本模板引擎示例
原始源代码 s 展开后 
[[for(int i=3; i<7; i++) {]]
  v_mov_b32 v[[=i]], v[[=i+4]] [[ } ]]
 
 → 
v_mov_b32 v3, v7
v_mov_b32 v4, v8
v_mov_b32 v5, v9
 
// Created [[=DateTime.Now]]
 → 
// Created 1/24/2015 8:12 PM


在同一个 cl::program 中混合 OpenCL 和 Asm4GCN 内核

___kernel 和 __asm4GCN 内核可以在同一个 OpenCL cl::program 中混合。最初的目标是让 OpenCL 内核具有内联功能，但这并不奏效——至少目前是这样。所以我决定采用单独的 __asm4GCN 和 __kernel 内核。这效果很好，因为这两种类型的内核都可以与流一起合并到 cl::program 中。注意：OpenCLwithGCN 的当前版本仅支持单个 __asm4GCN。（参见 限制）

使用 OpenCLwithGCN

为了展示如何使用 OpenCLwithGCN，我将逐步介绍项目中的 Example1.cs。

第一步是添加 OpenCLwithGCN.exe 和 NOpenCL.dll 作为项目引用。即使 OpenCLwithGCN 是一个可执行文件，它仍然可以作为引用添加。添加引用后，我们就可以开始编写代码了。

我们首先以几种方式之一创建一些 GCN 汇编。第一种方法是直接开始编写汇编。除非程序员确切地知道自己在做什么，否则这很困难，因为很难确定 参数将使用哪些寄存器。（参见 预加载寄存器值）另一种方法是编写一个简单的 OpenCL 内核，然后反编译它。反编译成汇编后，我们可以将汇编复制粘贴到我们的程序中并根据需要进行修改。

在此示例中，我将使用 example1.cs 中的 myOpenClFunc(...) 内核的反编译汇编（见下文）。example1.cs 中的 myOpenClFunc(...) 仅供参考。一个 cl::program 可以包含普通 OpenCL 内核和 asm4GCN 内核的混合。为了生成汇编，我使用了 AMD 的 CodeXL。

string source = @"
    __asm4GCN myAsmFunc ( float*, float* )
    {
      #define _32Float_    0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT]
      s_buffer_load_dword  s0, s[4:7], 0x04
      s_buffer_load_dword  s1, s[4:7], 0x18
      s_waitcnt            lgkmcnt(0)
      s_min_u32            s0, s0, 0x0000ffff
      s_buffer_load_dword  s4, s[8:11], 0x00
      v_mov_b32            v1, s0
      v_mul_i32_i24        v1, s12, v1 
      v_add_i32            v0, vcc, v0, v1
      v_add_i32            v0, vcc, s1, v0
      v_lshlrev_b32        v0, 2, v0
      s_load_dwordx4       s[12:15], s[2:3], 0x60 
      s_waitcnt            lgkmcnt(0)
      v_add_i32            v1, vcc, s4, v0
      tbuffer_load_format_x v1, v1, s[12:15], _32Float_
      s_buffer_load_dword  s0, s[8:11], 0x04
      s_load_dwordx4       s[4:7], s[2:3], 0x68
      s_waitcnt            lgkmcnt(0) 
      v_add_i32            v0, vcc, s0, v0
      s_waitcnt            vmcnt(0) 
      v_add_f32            v1, v1, v1
      tbuffer_store_format_x v1, v0, s[4:7], _32Float_ 
      s_endpgm 
    };

     __kernel void myOpenClFunc ( __global float* cl_input, __global float* cl_output )
     {
       size_t i = get_global_id(0);
       cl_output[i] = cl_input[i] + cl_input[i];
     }; ";

在下一步中，我们初始化 OpenClWithGCN 并编译 GCN。另外请注意，我们为了稍后使用而获取了 OpenClWithGCN 创建的默认环境。

OpenClWithGCN gprog = new OpenClWithGCN();
OpenClEnvironment env = gprog.env;
bool success = gprog.GcnCompile(source, out log);

此时，我们几乎完成了 OpenClWithGCN 的工作。在本例的其余部分，我将使用 NOpenCL 来设置 OpenCL 变量然后执行内核。我更喜欢 NOpenCL，因为它的包装器几乎与 C 语言的 OpenCL 一一对应。此外，本例的其余部分使用了 Derek Gerstmann (UWA) 编写的修改后的 OpenCL 示例。我试图采用这个例子，但我觉得我没有做到尽善尽美。

接下来，我们从修改后的程序创建一个 Kernel。

Kernel kernel = env.program.CreateKernel("myAsmFunc");

分配 cl_input 和 cl_output 设备缓冲区，并用随机数据填充它们……

Mem cl_input = env.context.CreateBuffer(MemoryFlags.ReadOnly, dataSz);
Mem cl_output = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);

// create some random data for testing
var random = new Random();
const int count = 1024 * 1024;
const int dataSz = count * sizeof(float);
float[] data = (from i in Enumerable.Range(0, count)
               select (float)random.NextDouble()).ToArray();

// Copy our host buffer of random values to the cl_input device buffer...
env.cmdQueue.EnqueueWriteBuffer(cl_input, true, 0, dataSz, data);

现在，让我们将缓冲区添加到内核中，并将其加入队列以执行

kernel.Arguments[0].SetValue(cl_input);
kernel.Arguments[1].SetValue(cl_output);

将内核加入队列并运行。对于这个内核，我们将使用 256 的工作组大小。

env.cmdQueue.EnqueueNDRangeKernel(kernel, count, 256);

强制命令队列处理，等待所有命令完成

env.cmdQueue.Finish();

读回结果

float[] results = new float[count];
env.cmdQueue.EnqueueReadBufferAndWait(cl_output, results, dataSz);

就是这样。现在只需验证并打印结果。

OpenCLwithGCN 项目文件

以下是此子项目的源代码文件及其描述的列表。

  • OpenClEnvironment.cs (97) - 此文件包含与保存 OpenCL 环境相关的类。
  • OpenClWithGCN.cs (615) - 此文件包含核心功能。
  • TextTemplate.cs (105) - 此文件包含一个单一的静态 Expand() 函数。它负责所有文本模板转换。

Asm4GcnGUI Windows 界面子项目

Asm4GcnGUI 为使用 GCN 汇编提供了一个图形用户界面 (GUI)。它比简单的记事本更容易使用，因为它具有语法高亮、代码补全和一些其他附加功能。此应用程序可用于进行一些快速内核测试、玩耍或用于教育目的。GUI 还充当如何实现 OpenCLwithGCN 的示例。启动时，它默认为一个简单的示例，可用作新项目的起点。该界面有三个窗口：C# 主机代码、内核代码和底部的编译器输出窗口。

在主机窗口中，您会找到在 CPU 上运行的 C# 代码。主机窗口包含构成正常 C# 应用程序的所有内容，因此可以将其放入 C# 项目中。要在 C# 中使用 OpenCL，需要一个包装器。该包装器将 C 语言的 OpenCL DLL 函数包装成 C# 兼容函数。我选择使用 NOpenCL——一个编写得很好的包装器。

下一个窗口是 GCN 汇编代码窗口。这是“C# Host Code”旁边的选项卡。通常，__asm4GCN 内核会与其余主机代码放在一起，但是 GUI 将其分开，以便为 GCN 汇编实现语法高亮。在后台，C# 主机和 GCN 汇编窗口中的文本会合并，然后一起编译。GCN 汇编窗口中的文本将被简单地放入一个 .CS 文件中，格式为 namespace GCN_NS { static class Code{ public const string DevCode = [GCN assembly Text here]}}”。这允许 GCN 汇编文本与主程序代码合并。

最后一个窗口是编译器输出。这显示了在编译过程中生成的任何错误、警告或消息。这里显示两组消息。首先是 Asm4GCN 汇编器消息，然后是任何 C# 主机代码错误。程序的输出不在此窗口中显示，它显示在单独的控制台窗口中。

GUI 界面的特点

  • 直接在 Visual Studio 中运行 – 差不多吧！我知道您可能不想长时间使用小型 IDE，因此在 Asm4GCN 中，保存文件时，它们将以可以直接在 Visual Studio 中打开的格式保存。您可能在想，GCN 汇编呢？它被包装在自己的命名空间和字符串常量中。唯一真正需要做的步骤是添加引用。如果文件结构没有过多修改，它仍然可以在 Asm4GCN 中重新打开。
  • 语法高亮 – 对于此项目，我使用了 Pavel Torgashov 的 Fast Colored Text box。它是一个很棒的控件，可以自定义以适用于 GCN 汇编等自定义语言。它提供语法高亮、代码折叠、热键支持、括号高亮、撤销、打印支持，等等。语法高亮有助于 可读性 ，也有助于处理错误。如果关键字输入错误，它将不会高亮显示。
  • 代码补全 – GCN 汇编 窗口支持代码补全。这有助于记住您可能正在寻找的命令。再次感谢 Pavel Torgashov 的这个插件。
  • 单独的 C# 和 GCN 汇编代码窗口 – 如前所述，这允许特定语言的语法高亮。它还将代码以逻辑方式分开：CPU 主机代码和 GPU 代码。

GPU 汇编编程通用话题

本节更多地与 GPU 汇编编程相关，而不是与 asm4GCN 相关。

GPU 汇编编程的优缺点

直接用 GCN 汇编编程有利有弊。这些是我根据我过去的汇编编程 经验 得出的个人观点。

优点

  • 潜在的高效和快速代码——如果程序员知道自己在做什么，他们通常可以创建更快、内存占用更小的代码。我估计熟练的手写内核性能可以提高 2 倍到 4 倍。
  • 汇编可以利用更多的硬件指令和功能。有些指令只能通过汇编访问。此外，只能通过使用汇编来利用某些指令选项。最后，GPU 处理器还有一些只能通过汇编访问的功能，例如特殊寄存器。
  • 玩转 GPU 汇编有助于程序员理解 GPU 的内部工作原理。这在编写 Cuda 或 OpenCL 内核时会有所帮助。
  • 在汇编中编程时，人类程序员通常比高级编译器能更有效地将项目放入内存。人类比编译器更有创造力。=)

缺点

  • 低级语言，如 GCN 汇编或 nVidia 的 PTX，比 OpenCL 和 Cuda 等高级语言编写起来花费的时间更长。
  • 难以维护——跟踪汇编函数不像高级语言那样容易。理解在不破坏东西的情况下要更改什么可能需要很长时间。通常，小的 5 到 20 行代码不会有太大问题。
  • 可读性低——汇编本质上不是为了帮助高级可读性而设计的。它提供的唯一可读性 是允许二进制代码人类可读。汇编的重构和抽象很少，这使得它难以阅读和理解。
  • GPU 汇编可能会在新一代 GPU 上失效。因为它专为特定芯片设计，所以在新一代 GPU 上可能会失效。
  • 根据 GPU 汇编二进制文件加载方式的不同，它们可能会在新驱动程序版本中失效。OpenCLwithGCN 子项目对此更敏感，因为它需要找到一个虚拟二进制文件以便用 GCN 作为
  • OpenCL 内核可以在包括 AMD GCN、nVidia、x86 和 x64 在内的不同架构上运行。汇编内核被锁定在 AMD GCN 1.x 架构上。幸运的是，AMD GCN 架构的开发周期较长，因此 GCN 不会经常改变。
  • 汇编更容易出错——高级语言旨在最大程度地减少常见的编程错误。
  • 当为了性能而用汇编编程时，可能需要一些努力才能胜过高级编译器。仅仅因为它用汇编编写，并不意味着它会更快。

总的来说，根据个人经验，汇编最适合代码中较小的、对时间要求严格的部分。大量的汇编代码如果布局不当会很快变得复杂。对于关键部分，1-50 行汇编是可以的，但随着内核变得越来越复杂，编译器开始占据优势，因为它能更好地跟踪事物。

预加载寄存器值

启动 GCN 内核时，某些寄存器具有预加载的值。对于新用户来说，这是一个必需的列表。由于驱动程序决定了其中的一些值，因此它们可能在未来的驱动程序版本中发生更改。

寄存器 名称
s[2:3] UAV 表指针
s[2:3] +0x60 base_resource_const1(#T)
s[2:3] +0x68 base_resource_const2(#T)
s[4:7] Imm Const Buffer 0
s[4:7] +0x00 Grid Size
s[4:7] +0x04 Local Size
s[4:7] +0x18 Base Global ID
s[8:11] Imm Const Buffer 1
s[8:11]+0x00 param1 offset 
s[8:11]+0x04 param2 offset
s[8:11]+0x08 param3 offset
s12 Group ID
v0 Local ID

GCN 汇编代码编写技巧

  • 提前规划 – 也许先用一种更高级的语言编写代码——这几乎就像伪代码。用一种更高级的语言编写函数可以完成几件事：
    • 它有助于理清细节，帮助您完全理解您正在编写的内容。对我来说，有时我认为我知道需要什么，结果却发现我没有完全理解问题。或者，我发现该函数甚至无法正常工作。在汇编中弄清楚这一点不是一个好主意。
    • 先用 OpenCL/Cuda 编写 GPU 函数可以提供一个大纲供参考。
    • 它能让思维做好准备。我经常发现第二次编写函数会导致函数更简洁、更精确。
  • 了解 GCN ISA 手册 - 查阅 AMD GCN ISA 编程手册，了解可用的指令，这些指令可能有助于完成内核的各项工作。
  • 展开代码以避免跳转 – 跳转会消耗额外的计算周期，而这些周期通常可以避免。
  • 限制内核大小以适合缓存 - 尽量让您的内核足够小，以适合计算单元的共享指令缓存（目前为 32kb）。（4000-8000 条指令）我从 Realhet 那里学到了这一点。
  • 限制寄存器使用 – 使用的寄存器越少，可以运行的延迟隐藏线程就越多。我发现的一个节省寄存器使用的小技巧是将指令在代码中向上或向下移动以降低总体寄存器使用量。例如：假设在第 10 行，给 A 和 B 赋值，然后在第 20 行它们首次使用“C=A+B”，最后 C 在第 30 行使用。寄存器使用量将是 10-20 行的 2 个变量，然后是 20-30 行的 1 个变量。总寄存器使用量可以写为（2 个变量 * (20-10)）+（1 个变量 * (30-20)）= 30。为了节省寄存器，我们可以将“C=A+B”移到第 11 行。这最大程度地减少了使用两个寄存器的时间。然后我们的使用量将是（2 个变量 * (11-10)）+（1 个变量 * (30-11)）= 21。基本上，我们只是从第 11-20 行释放了一个寄存器。这个技巧仅适用于汇编，因为 OpenCL 和 Cuda 足够智能，可以为您完成。
  • 随便玩玩，享受 GCN 汇编的乐趣。练习越多，您就会越好。
  • 阅读在线博客、帖子和文章（例如这篇）=)

局限性

  • 不支持第三代（火山岛）GPU。[Radeon R9 280、Fury 和 Nano]
  • 不适用于某些驱动程序版本（请参阅 系统要求）
  • 仅兼容具有 GCN 技术的 AMD 显示适配器（请参阅 系统要求）

未来愿望清单

  • 修正同一个变量名无法重用的问题。
  • GCN 第三代支持
  • 其他 OpenCL 2.0 支持
  • 增加对其他 GPU 驱动程序版本的兼容性。
  • 内联汇编 – 这是最初的目标，但生成虚拟内核存在问题。
  • 其他友好、易读的汇编语句。而不是像 v_mul_i32_i24 varA, varB, varC 这样的语句，也许可以像 varA = VarB * VarC. 那样。为此，我们需要知道变量的类型来选择正确的指令。这部分已经准备就绪。
  • 使用变量类型信息（I、U、F、B）在与不匹配的指令一起使用时发出警告。例如，使用 v4f 配合 v_mul_i32_i24

视频

我上传了至少一个教程视频。请注意，视频的某些部分可能已过时。更新视频不像更新文档那样容易。

https://www.youtube.com/results?search_query=asm4gcn

系统要求

  • 支持 GCN 1.0 和 1.1（Gen2）技术的 AMD 显卡。不支持第三代（火山岛）GPU，如 Radeon R9 280、Fury 和 Nano。
  • 支持 AMD 显卡驱动版本 13.251、14.501、15.200、15.201、15.300 和 16.150。其他版本可能无法工作。

历史

  • 2015 年 2 月 16 日 - 首次公开发布
  • 2015 年 3 月 1 日 - 一般修复和更改
    • 修复：变量总是使用寄存器 0
    • 修复：移除了单个 __asm4GCN 块限制 - 现在一个程序中可以有多个 __asm4GCN 内核。
    • 更改：参数名称已删除 - 由于参数名称未使用，保留它们可能会令人困惑。函数头现在采用以下形式：__asm4GCN myAsmAddFunc (float*,float*){...}
    • 更改：合并了 #ref 命令到普通变量声明中。由于 #ref 命令与普通变量声明几乎相同，只是指定了寄存器，因此最好将它们合并。这样更清晰，也不那么令人困惑。引用格式不再是  #ref s8u myVar s[2:3] ，而是 s8u myVar s[2:3]。
    • 改进：扩展了自动完成框 - 它现在更适合代码。
    • 改进：清理了示例代码。
    • 改进：语法高亮 - 现在它高亮显示标签、寄存器和定义。它还高亮显示匹配的单词。
    • 移除：移除了自动编译跳过功能。如果没有代码窗口中的更改，此功能将跳过重新编译。它被移除是因为它增加了代码的复杂性，而且由于编译过程本身非常快，所以几乎没有性能优势。
    • 添加：ren 命令 - 添加了一个重命名命令。这允许在变量的使用发生变化时重命名变量。
    • 更新：更新了本文档
  • 2015 年 4 月 22 日 - 发布于 GitHub
  • 2015 年 7 月 18 日 - 修复了自动完成和分支
    • 更改：OpenCL 包装器已切换为使用 NOpenCL ，由 Tunnel Vision Laboratories 开发。这是一个很棒的、编写精良的包装器。
    • 添加：变量索引。（例如 myVar[1] 将访问 myVar 中的第二个寄存器）
    • 添加：VINTRP 编码指令
    • 修复：修复了 SAPP 编码中的错误。跳转无法正常工作。
    • 修复：GCN 选项卡上的自动完成不起作用 - 已修复。
    • 更新：更新了本文档
  • 2015 年 8 月 2 日 - 处理了变量
    • 新功能：变量在使用它们的最后一行上自动释放。
    • 新功能：跳转现在可以使用在任何语句之前，而不仅仅是行。
    • 重构：添加了变量类并将变量功能移入变量类
    • 重构：重新组织了跳转功能的实现方式
    • 重构：以前所有操作都是单次通过，现在是两次通过
      • 1) 第一遍 - 读取所有语句并记录变量的位置
      • 2) 处理自动变量释放
      • 3) 处理变量的寄存器分配
      • 4) 第二遍 - 将语句转换为二进制
    • 移除：移除了 'ren' 函数，因为它产生了难看的代码。通过使用指定寄存器的变量声明可以达到相同的效果。
  • 2015 年 8 月 9 日 - 一般更新
    • 新功能：内联变量声明（例如 v_mov_b32 v4u myNewVar, anyVar）
    • 新功能：现在可以在一行中添加多个混合语句的标签
    • 新增：现在可以在同一条指令中重用已释放的变量寄存器。
    • 新增：以“/”结尾的行将追加下一行。#define 和语句可以跨越多行。
    • 新增：清理了初始代码，并添加了一个 #define(...) 来方便查看任何 S 或 V 变量。
    • 新增：为 Visual Studio 用户添加了 Ctrl-Y 作为重做操作符。
    • 新增：附加的变量警告检查（例如，当变量从未使用或仅使用一次时）。
    • 新增：声明变量时，可以将现有变量（带有选项索引）指定为要重用的寄存器。
      重用。
    • 更改：空格不能再用于分隔操作数。只能使用逗号。这不适用于末尾的指令选项参数。
    • 更改：现在先处理 #define，然后再处理标签。
    • 更改：#define 现在以反向顺序处理，因此 #define 替换可以包含之前的 #define 替换。 
    • 修复：多寄存器变量未始终正确对齐。
    • 修复：修复了变量索引问题。 
    • 修复：GCN 选项卡上的语法颜色高亮问题，并非所有文本都能始终正确高亮。
    • 修复：对指令编码器进行了一些小的修复和调整。
    • 新增：添加了单元测试，可兼作示例。
    • 新增：添加了示例：一个快速的波前求和规约，使用 18 条指令且不使用共享内存。

  • 2015 年 11 月

    • 新增：添加了友好的语法功能，以便可以更易读的格式输入 asm 语句。例如：“v_add_i32 localSizeIdx, vcc, laneId, vLocalSize”现在可以添加为“localSizeIdx = laneId + vLocalSize”。这更易于阅读。仅支持 +, -, *, >>, <<。
    • 重构：将部分正则表达式提取到单独的文件中。未来将对此进行扩展。
  • 2016 年 1 月 - 4 月 
    • 更新；添加了内核导出的提示。
    • 更新：将驱动程序版本 15.300 和 16.150 添加到可接受列表中；代码清理。
    • 修复：当存在多个 AMD GCN GPU 时，Asm4GCN 会失败。现在可以工作，但 Asm4GCN 在没有调整的情况下无法与两个 GPU 一起工作。
    • 更新：添加了对更新驱动程序版本的支持；为 '+' 的友好转换器添加了“vcc”。
    • 修复：修正了 DS_WRITE_SRC2_B32=141 和 DS_WRITE_SRC2_B64=205 的操作码（感谢 Mateusz Szpakowski 的发现和更正）。

其他 GCN 汇编器

cmingcnasm - cmingcnasm 是一个 GCN 最小汇编器，用 C 语言编写，可在 Linux 平台上与 GCC 一起使用。由 Sylvain Bertrand 创建。我没有试用过这个库，但我认为它可能很快，因为它是用 C 语言编写的。链接： GitHub,  GoogleCode

gcnasm - 由 Daniel Bali 创建，是一款开源的 GCN 汇编器，用 C 语言编写。代码使用低级 C 语言高效编写，因此速度极快。它在 Linux 中使用 GCC 编译器，但由于其标准的 C 包含文件，可以轻松修改以在 Windows 环境中工作。许多关于如何构建 Asm4GCN 的想法和代码参考都来自 gcnasm。 链接： GitHub,  OpenWall

HetPas Assembler (Windows)- 由 Realhet 编写，在这一点上，HetPas 是唯一可直接在 Windows 上运行的 GCN 汇编器。我使用 HetPas 的汇编器已有几年了。它非常适合创建二进制内核或仅是玩弄 GCN 汇编。该程序使用类似 Pascal 的语言为主机代码实现了功能齐全且丰富的 GCN 汇编器。它需要关闭 Windows 中的“数据执行保护”，并且有如何执行此操作的说明。HetPas 已更新多次，并且随着每次发布的增加，功能也越来越丰富。最近还添加了变量。 

特别感谢....

AMD 提供他们的 GCN ISA 手册。我花费了无数小时查阅 ISA 文档并参考他们的数据表。 

Daniel Bali, 他构建了一个优秀的开源 GCN 汇编器。这是我的第一个编译器，所以我一直在寻找如何开始的思路。Daniel 的项目给了我如何着手构建这个项目的灵感。我从他的项目中学习了一些构建汇编器的重要概念。

Derek Gerstmann 提供了他易于遵循且完整的 OpenCL 示例。Derek 的 OpenCL 示例是 Asm4GcnGUI 中默认打开的示例。它已被转换为 C# 以用于 NOpenCL。

Pavel Torgashov 提供了 FastColoredTextBox 编辑器 和 自动完成菜单。这些控件为 GUI 界面增加了语法高亮和代码补全功能。 

Realhet, 他为 Windows 构建了一个功能齐全且功能丰富的汇编器，名为 HetPas。我的许多 GCN 汇编技能都来自于玩弄 Realhet 的汇编器以及阅读他的帖子。我一直是 Realhet 的粉丝。他在 AMD 论坛 和他的 WordPress 网站上发表了许多富有洞察力的帖子。

Tunnel Vision Laboratories 提供了非常出色的 OpenCL 包装器，名为 NOpenCL。它由 Sam Harwell 编写。

© . All rights reserved.