@llvm.amdgcn.workitem.id.x()引发的一些前后端的调研-CSDN博客

本文链接：https://blog.csdn.net/u012276729/article/details/136641884

记录资料：

User Guide for AMDGPU Backend — LLVM 5 documentation

intrinsic函数会执行lowerintrinsics pass

llvm-project-main/llvm/lib/CodeGen/IntrinsicLowering.cpp

llvm-project-main/llvm/lib/Target/AMDGPU/AMDGPULowerIntrinsics.cpp

intrinsic函数的名称和语义可以是预先定义，也可以自定义，要求遵守特定的约定。在有些情况下，可能会调用库函数。例如，调用libc。

总的来说，这些Intrinsic函数代表了LLVM语言的一种扩展机制，当添加到语言中时，不要求改变LLVM的任何转化过程。对其它编译器，Intrinsic函数也称为内建函数。（摘自LLVM的Intrinsics函数及其实现 - 知乎）

在LLVM中，Intrinsic函数一般是在IR级代码优化时引入的，也就是由前端产生。也可以在程序代码中写Intrinsic函数，并通过前端直接发射。这些函数名的前缀一般是保留字“llvm.”。（plugin或者是内嵌汇编吗？）

#include<string.h>
int foo(void){
 char str[10] = "str";
 return 0;
}

define i32 @foo() #0 {
entry:
 %str = alloca [10 x i8], align 1
 %0 = bitcast [10 x i8]* %str to i8*
 call void @llvm.memcpy.p0i8.p0i8.i64(i8* %0, i8* getelementptr inbounds ([10 x i8]* @foo.str, i32 0, i32 0), i64 10, i32 1, i1 false)
 ret i32 0
}

llvm.memcpy就是clang输出的Intrinsic函数。如果LLVM没有定义llvm.memcpy，相应的内存操作LLVM IR代码就应该是一系列 "store constant into str[0..3]"内存访问指令，而这些指令通常都是极耗时的。LLVM后端可将llvm.memcpy拆分为一系列高效机器指令，也可以映射为一条特定的机器指令，直接调用硬件的内存操作功能。（对llvm ir中不高效的指令打包，替换，是一种优化）

libclc中定义的头文件：

libclc/generic/include/clc/workitem/get_work_dim.h

_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim);

#include <clc/clc.h>

_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
  return get_num_groups(dim)*get_local_size(dim);
}

当你编写OpenCL内核并使用get_work_dim()这样的内置函数时，你的代码会包含一些OpenCL的头文件（比如opencl.h），这些头文件会间接地引用到libclc中的相关实现或声明。在编译时，编译器会查找这些头文件，并解析其中的函数声明和定义，以确保你的代码能够正确地编译成OpenCL可执行内核。（这些函数的隐式参数如何传递到内存上？）

虽然libclc提供了这些函数的实现或声明，但实际的函数执行是在OpenCL运行时环境中完成的，而不是由libclc本身提供的库函数。因此，这个头文件的主要目的是帮助编译器在编译阶段理解你的OpenCL代码，而不是提供运行时的库函数实现。（运行时环境提供的？）我的理解：应该说opencl编译器都是运行时编译的，所以所谓的运行时环境提供函数实现，我认为还是intrinsic函数，并不是什么实现方式。

补充相关知识1：

运行时编译通常不是指动态编译，尽管这两个概念在某些上下文中可能被混淆。为了更准确地解释这两个概念，我们首先要明确它们各自的含义。

运行时编译（Runtime Compilation）：
在OpenCL的上下文中，运行时编译通常指的是在程序运行时，将OpenCL内核代码编译成可以在特定设备上执行的机器代码。这通常是通过调用OpenCL API中的函数（如clBuildProgram）来完成的。这个编译过程发生在程序执行期间，而不是在程序编译时。它允许程序在运行时根据设备的能力和其他条件来优化内核代码的执行。

动态编译（Dynamic Compilation）：
动态编译通常指的是在程序执行期间，根据需要动态地生成和编译代码。这可以是为了实现某种形式的代码生成、优化或适应性。动态编译与静态编译（即在程序编译时将所有代码编译成机器代码）相对。动态编译可以提供更多的灵活性和适应性，但也可能会带来额外的性能开销。

在OpenCL的上下文中，运行时编译可以看作是一种动态编译的形式，因为它允许在程序运行时动态地生成和编译内核代码。然而，并不是所有运行时编译都是动态编译。例如，某些运行时系统可能允许在程序启动前预先编译内核代码，并在运行时直接加载执行，这就不属于动态编译的范畴。

补充相关知识2：

运行时编译和静态编译在多个方面存在显著的不同。以下是它们之间的主要区别：

编译时机：
- 静态编译：在编译时将程序的源代码和所有依赖的库文件一起编译成一个可执行文件。这意味着在编译阶段，所有的代码和依赖都已经被确定和打包。
- 运行时编译：在程序运行时动态地生成和编译代码，或者根据需要将特定的代码段编译成机器代码。这意味着编译过程发生在程序执行期间。
执行文件大小：
- 静态编译：由于将所有依赖的库文件都编译进了可执行文件，因此生成的可执行文件通常较大。
- 运行时编译：可执行文件本身可能较小，因为它不包含所有的依赖库。这些库可能在运行时动态加载。
性能：
- 静态编译：由于所有的依赖都已经在编译时确定，并在可执行文件中，因此在运行时不需要查找和加载依赖库，这有助于提高程序的启动速度和执行效率。
- 运行时编译：在运行时可能需要查找和加载依赖的库文件，这可能会带来一定的性能损失，尤其是在首次加载时。然而，它也提供了更大的灵活性，允许程序在运行时根据需要进行优化或适应不同的环境。
灵活性：
- 静态编译：一旦编译完成，可执行文件就固定了，不容易进行后续的修改或优化。
- 运行时编译：由于编译发生在运行时，因此可以更容易地实现代码的动态生成、优化或适应不同的运行条件。
依赖管理：
- 静态编译：减少了运行时对外部库的依赖，使得部署更为简单，但也可能导致可执行文件体积较大。
- 运行时编译：需要确保在运行时所有必要的库都是可用的，这增加了部署的复杂性，但也可能使得程序更为轻量级和模块化。

那么回到libclc，libclc不是OpenCL运行时的直接依赖库呢？

libclc是一个开源项目，提供了OpenCL C语言库的C语言实现，它主要用于帮助实现OpenCL编译器或运行时环境的一部分功能。然而，OpenCL运行时环境本身并不直接依赖于libclc。

OpenCL运行时环境是负责执行OpenCL内核代码并在设备上管理资源的软件层。它提供了与设备通信、内存管理、内核编译和执行（cuda也是运行时环境编译吗？答案在后面）等功能。OpenCL运行时环境通常由设备供应商提供，并且与特定的硬件平台紧密相关。

虽然libclc提供了OpenCL内置函数的实现，但OpenCL运行时环境可以使用其他方式来实现这些内置函数，而不一定依赖于libclc。实际上，OpenCL运行时环境可能使用自己的内部实现或与其他库进行集成，以提供这些内置函数的功能。

因此，虽然libclc可以为OpenCL编译器和运行时环境的开发提供帮助，但它并不是OpenCL运行时的直接依赖库。OpenCL运行时环境通常具有自己的实现方式，并且可以根据需要进行定制和优化。（那么opencl运行时的自己实现函数的方式是？）

举一些opencl标准的内置函数：

OpenCL运行时环境确实实现了许多内置函数（也称为内建函数或intrinsics），这些函数在OpenCL C语言编程中可以直接调用，无需额外的库链接。这些内置函数是OpenCL规范的一部分，由OpenCL运行时环境提供实现，以确保它们能在不同的OpenCL设备上正确、高效地执行。

以下是一些OpenCL运行时环境通常实现的内置函数的例子：

内存访问函数：
- get_global_id(uint dimindx): 获取当前工作项的全局唯一ID。
- get_global_size(uint dimindx): 获取指定维度上的全局工作项数量。
- get_local_id(uint dimindx): 获取当前工作项在其工作组内的局部ID。
- get_local_size(uint dimindx): 获取指定维度上的工作组大小。
- get_group_id(uint dimindx): 获取当前工作组的全局ID。
- get_num_groups(uint dimindx): 获取指定维度上的工作组数量。
数学函数：
- sin(float x): 正弦函数。
- cos(float x): 余弦函数。
- tan(float x): 正切函数。
- sqrt(float x): 平方根函数。
- pow(float x, float y): 幂函数。
- exp(float x): 指数函数。
- log(float x): 自然对数函数。
- floor(float x): 向下取整函数。
- ceil(float x): 向上取整函数。
- 以及其他许多常用的数学函数。
原子操作函数：
- atomic_add(volatile __global *p, T val): 原子地给全局内存位置加值。
- atomic_sub(volatile __global *p, T val): 原子地从全局内存位置减值。
- atomic_min(volatile __global *p, T val): 原子地设置全局内存位置为当前值和给定值中的较小者。
- atomic_max(volatile __global *p, T val): 原子地设置全局内存位置为当前值和给定值中的较大者。
- atomic_xor(volatile __global *p, T val): 原子地对全局内存位置进行异或操作。
- 以及其他用于并发访问和修改全局内存的原子操作。
图像操作函数：
- 用于处理OpenCL图像对象的函数，如读取、写入和采样图像像素。
事件同步函数：
- wait_for_events(cl_uint num_events, const cl_event *event_list): 等待一个或多个事件完成。
- get_event_info(cl_event event, cl_event_info event_info, size_t param_value_size, void *param_value, size_t *param_value_size_ret): 查询事件的信息。

这些内置函数在OpenCL内核代码中可以直接调用，而无需额外的库链接或显式地实现它们。OpenCL运行时环境负责为这些函数提供高效、可移植的实现，以适应不同的硬件平台和架构。

应该说这些内建函数的函数名都是在libclc中申明的，可能有的在libclc中已经实现了吧，这个还么调研过，但是这里的函数名是程序员可见的，在llvm编译器中都是由后端的intrinsiclower pass 转换到成后端对应的的intrinsic函数。下一步，在selectDAG的时候，这些被intrinsic函数会转换成相关的硬件SDValue，也就是硬件指令或者寄存器，内存地址等，继续做指令调度，寄存器分配。

这里想到之前调研的nvvm ir，所以想在这里在继续深入调研一下：

NVVM IR（NVIDIA Virtual Machine Intermediate Representation）是基于LLVM IR的，专门为表示GPU计算kernel（例如CUDA kernel）而设计的编译器IR。NVVM IR的引入使得NVIDIA的GPU编译器能够更有效地处理CUDA代码，进行转换和优化。
Clang编译CUDA代码生成的中间表达式是NVVM IR。（是吗？）NVVM IR（NVIDIA Virtual Machine Intermediate Representation）是NVIDIA为其GPU编译器提供的一种中间表示形式，用于在CUDA代码和其他GPU代码之间进行转换和优化。Clang是LLVM项目的一部分，它可以将CUDA代码转换为NVVM IR，然后由NVIDIA的GPU编译器进一步处理以生成可在GPU上执行的代码。

虽然LLVM IR是LLVM项目的一种通用中间表示形式，用于在各种编译器前端和后端之间进行转换，但在处理CUDA代码时，Clang选择使用NVVM IR而不是LLVM IR，因为NVVM IR更适合GPU代码的特殊需求。

尽管NVVM IR主要关注CUDA代码，但NVIDIA的GPU编译器工具链可能提供了一些机制，使得开发者能够使用其他编程语言（如C++、Fortran等）编写GPU代码，并将这些代码转换为NVVM IR，进而生成可在GPU上执行的代码。这种转换通常是通过前置编译器或特定的编程接口实现的，这些前置编译器或接口将其他编程语言的代码转换为CUDA代码或类似的中间表示形式，然后再进一步转换为NVVM IR。（这倒是和cuda兼容反过来了，tvm后端不知道是直接转换成cuda代码还是某种中间表示呢）

也就是说，NVVM IR主要用于优化CUDA代码，但NVIDIA的GPU编译器工具链可能提供了将其他编程语言转换为NVVM IR的机制，以实现跨语言的GPU编程。（好吧，哪些工具链呢？找出来瞧瞧）

如果跳过nvvm ir，其他的cuda编译方式是什么呢？

clang编译cuda项目生成的设备代码对应的llvm ir如何转换成ptx，答案是nvptx。

NVPTX（NVIDIA PTX）的工具链主要包括以下几个组件：

Clang/LLVM：Clang 是一个基于 LLVM 的 C/C++/CUDA 编译器前端，它可以将源代码编译成 LLVM 中间表示（IR）。对于 CUDA 设备代码，Clang 可以生成适用于 NVPTX 的 LLVM IR。LLVM 则提供了一组编译器和工具链技术的集合，包括将 LLVM IR 转换为 PTX 代码的功能。
NVCC（NVIDIA CUDA 编译器）：虽然 NVCC 不是开源的，但它是 NVIDIA 官方提供的 CUDA 编译器，用于将 CUDA 源代码编译成 GPU 可执行的二进制代码或 PTX 代码。NVCC 在 CUDA 开发中扮演着重要角色，但它与 NVPTX 的开源工具链不完全相同。然而，NVCC 可以与 Clang/LLVM 结合使用，以利用 LLVM 的优化和代码生成能力。
PTXAS（PTX Assembler）：PTXAS 是 NVIDIA 提供的一个工具，用于将 PTX 代码汇编成 GPU 二进制代码（即 CUBIN）。它是 CUDA 工具链的一部分，通常与 NVCC 一起使用。然而，PTXAS 本身不是开源的。
CUDA Runtime API 和 CUDA Driver API：这些 API 提供了与 CUDA 设备进行交互的接口和功能。CUDA Runtime API 提供了一组高级别的函数，用于管理 CUDA 设备、内存分配、数据传输和内核启动等任务。CUDA Driver API 则提供了更低级别的接口，允许更精细的控制和更高的灵活性。这些 API 通常是闭源的，但它们是 CUDA 开发中不可或缺的一部分。

尽管 Clang/LLVM 提供了生成 NVPTX 代码的能力，但完整的 CUDA 开发体验仍然需要依赖 NVIDIA 提供的闭源组件，如 CUDA 驱动程序、运行时库和 NVCC 编译器。这些组件提供了与 GPU 硬件进行交互所需的底层接口和功能。

另外，开源社区也开发了一些与 CUDA 和 NVPTX 相关的工具和库，如 CUDA Toolkit 中的 CUDA Samples、cuDNN（CUDA Deep Neural Network library）等。这些工具和库可以帮助开发人员更高效地开发 CUDA 应用程序，但它们并不是 NVPTX 工具链的核心组成部分。

Clang 编译 CUDA 项目时，生成的设备代码对应的 LLVM IR 可以通过 llc转换成 PTX（Parallel Thread Execution）代码，这依赖于NVIDIA 的 nvptx 后端，也即是llvm中的nvptx target。PTX 也是一种中间表示形式，它会在 NVIDIA 的 GPU 驱动程序在运行时将代码编译成特定 GPU 架构的机器代码。（这不又是intrinsic函数嘛，但是它又有自己的指令集，又有很多intrinsic函数，比如矩阵乘法，如果是这样的话，它和llvm ir确实也是一样的形式，都有指令和函数。selectDAG的时候就会替换具体架构的指令SDValue，和intrinsic函数替换）

nvvm ir就到此为止吧，作为llvm ir的一个分支而已，现在有MLIR了，不如去学MLIR。不过两个都很相似，就当nvvm ir就算是llvm ir的方言吧。

如果想要尝试将 Clang/LLVM 与 NVCC 结合使用，可以考虑以下步骤：

然而，需要注意的是，这种结合使用的方式并不是 NVIDIA 官方推荐的标准工作流程。通常，CUDA 开发者会直接使用 NVCC 来编译 CUDA 源代码，因为它提供了针对 NVIDIA GPU 的优化和特性支持。Clang/LLVM 的支持主要用于那些对编译器技术有深入了解并希望进行更高级定制或研究的开发者。

编译 CUDA 代码到 LLVM IR：
使用 Clang 编译 CUDA 源文件，生成 LLVM IR。这通常涉及使用特定的编译选项来指示 Clang 生成 IR 而不是直接生成机器代码。例如，可以使用 -emit-llvm 选项。

clang -cc1 -triple nvptx64-nvidia-cuda -emit-llvm -o output.ll input.cu
请注意，上面的命令可能需要根据您的 Clang 版本和 CUDA 安装进行调整。特别是，-triple 选项的值可能需要与您的目标 GPU 架构相匹配。
（可选）优化 LLVM IR：
你可以使用 LLVM 的优化工具（如 opt）对生成的 LLVM IR 进行优化。这些优化可能包括死代码消除、常量折叠、循环展开等。
将 LLVM IR 转换为 PTX：
一旦有了 LLVM IR，您可以使用 LLVM 的 llc 工具（LLVM 静态编译器）将其转换为 PTX。您需要确保使用 NVIDIA 提供的 nvptx 后端。

llc -march=nvptx64 output.ll -o output.ptx
同样，-march 选项的值应该与您的目标 GPU 架构相匹配。例如，对于基于 Turing 架构的 GPU，您可能需要使用 nvptx64 或更具体的架构标识符。
链接和加载 PTX：
生成的 PTX 代码随后可以被 CUDA 运行时库加载并编译成 GPU 机器代码。这通常是在 CUDA 应用程序运行时自动完成的，但也可以手动加载和编译 PTX。
手动使用 NVCC 进行链接和最终编译。尽管 LLVM IR 已经被转换为 PTX，但手动的话，仍然需要 NVCC 来执行最终的链接步骤和/或将 PTX 代码编译为 GPU 可执行的二进制代码（CUBIN），通常涉及使用 NVCC 的 -cubin 选项或其他相关选项。

Clang 对 CUDA 的支持可能因版本而异，并且可能不完全等同于使用 nvcc。因此，在将 Clang 用于生产环境的 CUDA 项目之前，确保测试并验证生成的代码的正确性和性能。

另外，NVIDIA 提供了 nvrtc（NVIDIA Runtime Compiler）库，它允许在运行时将 CUDA C/C++ 代码编译成 PTX 或 CUDA 二进制代码。这对于需要动态生成或编译内核的应用程序很有用。然而，这与使用 Clang 生成 LLVM IR 再转换为 PTX 的过程不同，它是直接在运行时编译源代码。（这个应该和opencl是一样的吧。）

那就看看nvptx

研究一下intrinsic函数：首先得看看selectionDAG模块

在LLVM（Low Level Virtual Machine）中，SelectionDAG（Selection Directed Acyclic Graph）是一个将LLVM中间表示（IR）降低为机器代码的关键组件。SelectionDAG实际上是一个特定类型的DAG（有向无环图），它在指令选择阶段表示程序的低级形式。下面是一个高级概述，说明了SelectionDAG如何处理IR中的函数：

函数分解：首先，LLVM IR中的每个函数都被视为一个独立的单元进行处理。SelectionDAG的构建通常是从函数的入口开始，并按照控制流的顺序逐步进行。
基本块和DAG构建：IR中的每个基本块（一组顺序执行、无分支的指令）都被转换为一个或多个SelectionDAG节点。这些节点代表了机器操作或更接近机器的抽象，并且它们的连接关系表示了数据和控制依赖关系。
指令选择：一旦构建了代表整个函数（或函数的一部分）的SelectionDAG，就可以开始进行指令选择。这个过程中，DAG中的每个节点都被映射到一个或多个目标机器指令。这些指令可能直接对应于目标处理器的操作，或者是稍后在代码生成阶段进一步细化的抽象指令。
合并和优化：SelectionDAG的构建和指令选择过程可能会交错进行，以允许中间阶段的优化。例如，公共子表达式消除（CSE）可以在SelectionDAG级别进行，以移除重复的计算。
控制流图（CFG）的处理：除了表示基本操作和数据流之外，SelectionDAG还需要能够处理IR中的控制流。这通常是通过在DAG中插入特殊的控制流节点（如条件分支和函数调用节点）来实现的。
调度和寄存器分配：在指令选择之后，SelectionDAG还需要进行指令调度和寄存器分配。这些步骤决定了生成的机器指令的执行顺序以及它们将使用哪些硬件寄存器。这些过程可能会根据目标处理器的特定微架构特性进行优化。
代码发射：最后，一旦指令被选择、调度并分配了寄存器，SelectionDAG就可以被“降低”或转换为最终的机器代码。这个过程涉及将DAG中的节点转换为目标处理器的机器指令，并将这些指令以适当的格式（如汇编语言或二进制机器代码）输出。

intrinsic的匹配基本都在这个阶段实现。

下图中的pass，应该是也是直接替换intrinsic函数的，不知道什么时候用啊？

bool LowerIntrinsics::DoLowering(Function &F, GCStrategy &S) {
  SmallVector<AllocaInst *, 32> Roots;

  bool MadeChange = false;
  for (BasicBlock &BB : F)
    for (BasicBlock::iterator II = BB.begin(), E = BB.end(); II != E;) {
      IntrinsicInst *CI = dyn_cast<IntrinsicInst>(II++);
      if (!CI)
        continue;

      Function *F = CI->getCalledFunction();
      switch (F->getIntrinsicID()) {
      default: break;
      case Intrinsic::gcwrite: {
        // Replace a write barrier with a simple store.
        Value *St = new StoreInst(CI->getArgOperand(0),
                                  CI->getArgOperand(2), CI);
        CI->replaceAllUsesWith(St);
        CI->eraseFromParent();
        MadeChange = true;
        break;
      }
      case Intrinsic::gcread: {
        // Replace a read barrier with a simple load.
        Value *Ld = new LoadInst(CI->getType(), CI->getArgOperand(1), "", CI);
        Ld->takeName(CI);
        CI->replaceAllUsesWith(Ld);
        CI->eraseFromParent();
        MadeChange = true;
        break;
      }
      case Intrinsic::gcroot: {
        // Initialize the GC root, but do not delete the intrinsic. The
        // backend needs the intrinsic to flag the stack slot.
        Roots.push_back(
            cast<AllocaInst>(CI->getArgOperand(0)->stripPointerCasts()));
        break;
      }
      }
    }

  if (Roots.size())
    MadeChange |= InsertRootInitializers(F, Roots);

  return MadeChange;
}

编译器LLVM-MLIR-Intrinics-llvm backend-instruction_llvm mlir-CSDN博客

https://www.cnblogs.com/wujianming-110117/p/16508427.html

User Guide for AMDGPU Backend — LLVM 5 documentation

User Guide for AMDGPU Backend — LLVM 19.0.0git documentation

[AMDGPU] CodeGen for GFX12 64-bit scalar add/sub (#75070) · llvm/llvm-project@8005ee6 · GitHub

AMD 驱动程序与支持 | AMD

AMD GPU指令集是开源的吗？ - 知乎

How to accelerate AI applications on RDNA 3 using WMMA - AMD GPUOpen

https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/gcn3-instruction-set-architecture.pdf

PTX ISA :: CUDA Toolkit Documentation

OpenCL学习资料整理 - 知乎

https://github.com/nvidia/cuda-samples