llvm 编译器

LLVM 学习记录：如何添加一个 builtin 函数

Posted by 叉叉敌 on May 21, 2024

学习过程，求勘误，多谢指出~

今天学习如何在后端添加一个 builtin 和对应的指令 mapping 起来？

整个过程预览：

  1. 源代码调用：在高级语言的源代码中使用特定于 AMDGPU 的内置函数，例如就这么写 __builtin_amdgcn_sin(x)。

  2. LLVM IR 生成：Clang 编译器在编译过程中将这些内置函数调用转换为 LLVM IR 中的内置函数，例如 %0 = call float @llvm.amdgcn.sin.f32(float %x)。

  3. 机器代码生成：LLVM 的后端（例如 AMDGPU 后端）读取 LLVM IR，并根据 IntrinsicsAMDGPU.td 文件中的定义，将这些内置函数映射到具体的机器指令，就是 amd 的 isa 指令。

1. 什么情况下会添加这个函数？

https://llvm.org/docs/ExtendingLLVM.html

在使用 LLVM 的过程中，您可能希望为您的研究项目或实验定制它。此时，您可能会意识到需要向 LLVM 添加一些内容，无论是新的基本类型、新的内部函数还是全新的指令。

一般情况时有新的 target 添加，或者是支持新的 target，那么才需要，不然似乎需要。

因为这个需要花费大量的时间，请在社区咨询和是不是已经支持，或者已经有人在做了。

2. 开始添加

向 LLVM 添加新的内部函数比添加新指令要容易得多。几乎所有 LLVM 扩展都应该作为内在函数开始，然后在必要时转换为指令。

首先，要知道 tablegen 的语法。

关于 tablegen, 可以看下这篇文章，太多了，死机是记不住的，怎么办？多在实际项目中应用，记不住的再来查字典。

https://blog.llvm.org/posts/2023-12-07-tools-for-learning-llvm-tablegen/

https://llvm.org/docs/TableGen/

https://llvm.org/docs/TableGen/ProgRef.html

1. 添加定义 builtin 和 intrinsic，为了高级语言可以调用

首先需要了解点，基础的知识。

builtin 和 intrinsic 有什么区别？

一开始也不懂，只是啥呀？看名字感觉都是内置函数，确实都是内置函数。

我理解的有点区别，就是 builtin 是高级语言级别的，而 intrinsic 是 IR 级别的。

在 llvm 里面他们的命名大概是这样子的。

  • builtin：__builtin_memcpy()、__builtin_popcount()
  • intrinsic: llvm.memcpy(), llvm.ctpop()

CGBuiltin.cpp 这个文件干嘛的？

CGBuiltin.cpp 文件是 LLVM 项目的一部分，主要负责生成 LLVM IR 代码以实现内置函数的编译，包括初始化内存、生成内置库函数调用以及处理特定于目标架构的内置函数。

这个文件里面主要关注 2 个函数：文件包含了初始化内存的函数 initializeAlloca，该函数根据不同的初始化策略（如未初始化、全零初始化或模式初始化）来初始化 alloca 分配的内存。还有就是对齐方式，这个主要取决于硬件。

getBuiltinLibFunction 函数负责将高级语言中的内置函数映射到 LLVM IR 的内置函数，这对于生成有效的目标代码至关重要。

举个例子：比如高级语言是__builtin_fabs，对应的 IR 可能就是 llvm.fabs。

可以看到很多枚举值。比如：case RISCV::BI__builtin_riscv_ntl_store:

这个枚举值是通过 TargetBuiltins.h 里面通过 namespace 里面的 include 来的。

/// AMDGPU builtins
  namespace AMDGPU {
  enum {
    LastTIBuiltin = clang::Builtin::FirstTSBuiltin - 1,
  #define BUILTIN(ID, TYPE, ATTRS) BI##ID,
  #include "clang/Basic/BuiltinsAMDGPU.def"
    LastTSBuiltin
  };
}

Builtins*.def 是干嘛的

来到clang/Basic/BuiltinsAMDGPU.def这个文件里面，就可以看到具体的定义 BUILTIN 定义了。

有必要对这个代码解释下，我看了之后，一头雾水。LastTIBuiltin 和 FirstTSBuiltin 是什么？

  • 第一个是 namespace 这个好理解；
  • enum 枚举值；匿名枚举值；
  • LastTIBuiltin = clang::Builtin::FirstTSBuiltin - 1,：我看了大概是 AMDGPU 内置函数的 ID 值比 clang 内置函数的 ID 值小 1。 LastTIBuiltin：这个枚举成员表示 LLVM/Clang 中最后一个 Target-Independent (TI) 内置函数的编号。Target-Independent 内置函数是不依赖于特定硬件架构的内置函数。LastTIBuiltin 的值是 clang::Builtin::FirstTSBuiltin - 1，其中 FirstTSBuiltin 表示第一个 Target-Specific (TS) 内置函数的编号。这样设置是为了确保 AMDGPU 的内置函数枚举值在数值上紧跟在 TI 内置函数之后。
  • #define BUILTIN(ID, TYPE, ATTRS) BI##ID, 定义了宏，包含 3 个参数；展开后，应该类似 BI1，BI2……
BUILTIN(__builtin_amdgcn_dispatch_ptr, "v*4", "nc")
BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "v*4", "nc")
BUILTIN(__builtin_amdgcn_implicitarg_ptr, "v*4", "nc")
BUILTIN(__builtin_amdgcn_queue_ptr, "v*4", "nc")

估计你和我一样，第一眼看到的饿时候，what？这些神秘的 v，n，c 是什么意思？没关系，看文件的顶部，有一个文件，builtins.def，跟过去，看一下里面的内容，下面是部分内容。

//https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/Builtins.def

// (result value, then each argument) as follows:
//  v -> void
//  b -> boolean
//  c -> char
//  s -> short
//  i -> int
//  h -> half (__fp16, OpenCL)
//  x -> half (_Float16)
//  y -> half (__bf16)
//  f -> float
//  d -> double
//  z -> size_t
//  w -> wchar_t
//  F -> constant CFString
//  G -> id
//  H -> SEL
//  M -> struct objc_super
//  a -> __builtin_va_list
//  A -> "reference" to __builtin_va_list

……

// Types may be prefixed with the following modifiers:
//  L   -> long (e.g. Li for 'long int', Ld for 'long double')
//  LL  -> long long (e.g. LLi for 'long long int', LLd for __float128)
//  LLL -> __int128_t (e.g. LLLi)
//  Z   -> int32_t (require a native 32-bit integer type on the target)
//  W   -> int64_t (require a native 64-bit integer type on the target)
//  N   -> 'int' size if target is LP64, 'L' otherwise.
//  O   -> long for OpenCL targets, long long otherwise.
//  S   -> signed
//  U   -> unsigned
//  I   -> Required to constant fold to an integer constant expression.
//

看了就知道是什么意思了吧，v 就是 void，没有返回值。如果是Si signed integer 就是有符号整型。

Intrinsics*.td 是干嘛的

最后这个枚举值需要定义在llvm/IR/Intrinsics*.td 这个里面，就是第一步我们定义的。

llvm/include/llvm/IR/Intrinsics*.td:

用 amdgpu 来举例：

https://github.com/llvm/llvm-project/blob/main/llvm/include/llvm/IR/IntrinsicsAMDGPU.td

下面是我为了举例自己构造的，用法不一定正确。

def int_amdgcn_rcp_legacy : ClangBuiltin<"__builtin_amdgcn_rcp_legacy">,
  DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_i32_ty],
  [IntrNoMem, IntrSpeculatable]
>;
  • IntrNoMem 这些就是属性，也是定义在 Intrinsics.td 里面的，有解释，比如这个的定义意思就是：IntrNoMem - The intrinsic does not access memory or have any other side effects. It may be CSE’d deleted if dead, etc.
  • llvm_i32ty 是定义的一个数据类型属性，llvm 是命名空间的一部分，用于避免与其他库或框架中的类型名冲突。i32 表示 32 位整数，_ty 表示这是一个类型（type）的定义。

其他就去这个上面的链接上看。

顺便提一下这个 def 是啥和用法。

  • def 是定义一个 record：def 的语法 
    Def       ::=  "def" [NameValue] RecordBody
NameValue ::=  Value (parsed in a special mode)

这里的 intrinsic def 用法这么定义的（这个定义我没有找到，为什么这么约束的？我猜应该是在 Intrinsic.td 里面）

def <name> <args> <ret> [<attrs>] {
  // Definition body
}
  • name>：正在定义的实体的名称，例如内部函数、指令或数据类型。
  • args>：实体的参数列表，通常由其数据类型表示。
  • ret>：实体的返回类型（如果适用）。
  • attrs>：实体的可选属性列表，提供附加信息。 {}：包含定义主体，其中包含有关实体的特定详细信息。

这里的ClangBuiltin<"__builtin_amdgcn_rcp_legacy">就是 NameValue，就是可以在高级语言直接调用的。 RecordBody 就是后面的DefaultAttrsIntrinsic<>.

这个 Clangbuiltin 和 DefaultAttrsIntrinsic 都定义在：https://github.com/llvm/llvm-project/blob/main/llvm/include/llvm/IR/Intrinsics.td

// Intrinsic with default attributes (disable_default_attributes = false).
class DefaultAttrsIntrinsic<list<LLVMType> ret_types,
                list<LLVMType> param_types = [],
                list<IntrinsicProperty> intr_properties = [],
                string name = "",
                list<SDNodeProperty> sd_properties = []>
                : Intrinsic<ret_types, param_types,
                            intr_properties, name,
                            sd_properties, /*disable_default_attributes*/ 0> {}

定义好了，高级语言可以条用的接口，那么接下来就是，后端的实现了。

2. 把 intrinsic 和指令关联起来

需要到 llvm 的前端把 builtin 函数转化为 intrinsic，这个是前端实现的，我不太熟悉，我用 AI 搜索下了，回答如下：


1. 语法分析

首先，LLVM 的前端分析器（如 Clang）会对源代码进行语法分析，将源代码转换为抽象语法树（AST）。AST 包含了代码的结构和语义信息，包括函数声明、函数调用、表达式等。

在 AST 中，intrinsic 函数会被表示为特殊的节点，其中包含了函数名、参数类型和返回值类型等信息。

2. 语义分析

接下来，LLVM 会进行语义分析，检查代码的语义是否正确。在这个阶段，LLVM 会检查 intrinsic 函数的调用是否合法，例如参数类型是否匹配、函数是否被定义等。

如果语义分析成功，LLVM 会将 AST 转换为 IR。

3. IR 生成

在 IR 生成阶段，LLVM 会为每个 intrinsic 函数生成对应的 IR 函数。IR 函数由一系列 LLVM IR 指令组成，这些指令描述了函数的具体操作

然后是 intrinsic 和指令的映射。

这类文件 llvm/lib/Target/AMDGPU/*Instructions.td 就是定义具体的指令的，

在这个文件里面 https://github.com/llvm/llvm-project/blob/main/llvm/lib/Target/AMDGPU/DSInstructions.td 可以看到关于 ds_consume 的多个定义。

def DS_CONSUME       : DS_0A_RET<"ds_consume">;
defm DS_CONSUME             : DS_Real_gfx6_gfx7_gfx10_gfx11_gfx12<0x03d>;
def DS_CONSUME_vi         : DS_Real_vi<0xbd, DS_CONSUME>;

看一个，就可以看到这个类 DS_0A_RET的定义。

class DS_0A_RET <string opName> : DS_Pseudo<opName,
  (outs getLdStRegisterOperand<VGPR_32>.ret:$vdst),
  (ins Offset:$offset, gds:$gds),
  " $vdst$offset$gds"> {

  let mayLoad = 1;
  let mayStore = 1;
  let has_addr = 0;
  let has_data0 = 0;
  let has_data1 = 0;
}

里面就是约束了定义时候的用法，其他就没有了，还有一些属性。

再来看另外一个，类名DS_Real_vi。用的是一个 64 位的寄存器，0~63，每一位代表的是什么，都写了。

class DS_Real_vi <bits<8> op, DS_Pseudo ps> :
  DS_Real <ps>,
  SIMCInstr <ps.PseudoInstr, SIEncodingFamily.VI> {
  let AssemblerPredicate = isGFX8GFX9;
  let DecoderNamespace = "GFX8";

  // encoding
  let Inst{7-0}   = !if(ps.has_offset0, offset0, 0);
  let Inst{15-8}  = !if(ps.has_offset1, offset1, 0);
  let Inst{16}    = !if(ps.has_gds, gds, ps.gdsValue);
  let Inst{24-17} = op;
  let Inst{25}    = acc;
  let Inst{31-26} = 0x36; // ds prefix
  let Inst{39-32} = !if(ps.has_addr, addr, !if(ps.has_gws_data0, data0{7-0}, 0));
  let Inst{47-40} = !if(ps.has_data0, data0{7-0}, 0);
  let Inst{55-48} = !if(ps.has_data1, data1{7-0}, 0);
  let Inst{63-56} = !if(ps.has_vdst, vdst{7-0}, 0);
}

我强行解读下，估计是小尾端编码：

  • 0~15 位，都是表示 offset
  • 16 位：是一个 gds 的值，估计是一个 bool 啥的
  • 17-24：是一个操作数，opcode，每个指令不一样，是唯一的
  • 25：acc 是加速？准确？不装了，确实不懂……
  • 32-55 都是 data，应该是 source 之类的 addr
  • 56~63 是 destination，

3. 用 amdgpu 的来举例

首先，需要 https://github.com/llvm/llvm-project/blob/main/llvm/include/llvm/IR/IntrinsicsAMDGPU.td 里面定义个 intrinsic :: amdgcn_div_fixuph 函数。必须是 int_ 开头。

def int_amdgcn_div_fixuph : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
  [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
  [IntrNoMem, IntrSpeculatable]
>;

然后，在 https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/BuiltinsAMDGPU.def 文件里面定义个 builtin 的函数。

TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts")

然后，在 https://github.com/llvm/llvm-project/blob/main/clang/lib/CodeGen/CGBuiltin.cpp 里面，然后对应的 intrinsic 函数，这里返回的Intrinsic::amdgcn_div_fixuph。

case AMDGPU::BI__builtin_amdgcn_div_fixuph:
    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixuph);

这样前端就可以调用了，但是在 amdgpu 上还没有映射具体的指令。

这个需要在 intrinsic 里面添加这个定义。

defm V_DIV_FIXUP_F16 : VOP3Inst <"v_div_fixup_f16", VOP3_Profile<VOP_F16_F16_F16_F16>, AMDGPUdiv_fixup>;

这个对于的指令就是v_div_fixup_f16。

所以高级语言调用__builtin_amdgcn_div_fixuph 然后根据参数，会具体映射到定义的指令上面：v_div_fixup_f16。

end

CATALOG
    FEATURED TAGS
    基础 Newsletter llvm 编译器 python 测试 LLVM 平假名 日语 片假名 总结 cuda TensorFlow 大数据 compiler 工具 投资 git github 书单 写作 后端 提升 英语 逻辑 面试 Maven docker gRPC gcc optimization redis 保险 增长 技术 智能计算系统 简历 算法 线程 计算机 读书 调试 进阶 通信 音频 debug grep linux vscode webrtc
    Read More