jc小小川+幻幻融hr-CSDN博客

原创 tvm的IR定义、计算类型、数据类型和infer shape

*Infer Shape** 是编译过程中的一个重要环节，用于推断出模型中所有张量的形状（即维度和大小），这对于后续的编译优化和代码生成至关重要。2. **动态shape推断**：对于形状不确定的张量（如依赖于运行时输入的张量），通过分析计算图的结构和运算符的属性，结合用户提供的形状约束或运行时信息，递归地推断出其可能的形状范围或具体形状。1. **静态shape推断**：对于具有明确形状信息的张量（如在模型定义时指定的形状或通过`relay.shape_of`获取的形状），直接记录其形状信息。

2024-04-17 16:55:12 478

原创 cuda算子模板移植到自研芯片

**硬件特性**：了解自研芯片的并行计算单元（如SIMD、MIMD、矢量单元等）、内存层次（如缓存、全局内存、本地内存等）、访存模型（如一致性模型、带宽、延迟等）以及任何特定的加速硬件（如张量核心、AI加速器等）。- **编程模型**：掌握自研芯片提供的编程接口、编程语言（如专有ISA的汇编、高级语言绑定、DSL等）和并行编程框架（如类似CUDA的并行编程模型，或者完全不同的模型如任务图、数据流等）。- **性能调优**：根据自研芯片的特性，对移植后的代码进行细致的性能分析和优化。

2024-04-17 16:45:41 832

原创基于数据流驱动的matrix core 和 vector core设计

3. **编译器优化**：编译器和自动性能调优工具可以识别矩阵和向量操作，并自动将它们映射到Matrix Core和Vector Core上执行，无需开发者手动优化。1. **硬件实现**：Matrix Core和Vector Core通常作为GPU内部的专用硬件模块实现，它们直接集成在GPU的流式多处理器（SM）中。这些操作通常可以通过并行加法和减法单元来实现。3. **数据路径优化**：为了最大化吞吐量，Vector Core的数据路径会被优化，以支持高带宽的数据传输和快速的寄存器间数据移动。

2024-04-17 16:36:58 331

原创 AI算子分析，包含以下一个方面：优先级，计算pattern，芯片需求

一些关键的算子，如卷积、激活函数和池化等，在深度学习模型中扮演着至关重要的角色，因此具有较高的优先级。例如，一些硬件平台可能更擅长处理连续的向量运算，而对于这种类型的硬件，优化卷积算子的计算pattern可能会带来显著的性能提升。对延迟敏感的算子主要指的是那些在深度学习或机器学习模型中，执行时间对整体性能有关键影响的算子。为了减少这些算子的延迟，可以采取多种优化策略，如使用更快的数学库、利用硬件加速（如GPU或TPU）、优化算法实现、减少不必要的计算（如剪枝和量化）以及并行计算等。

2024-04-17 16:10:26 301

原创 Dispatch架构

GPU（Graphics Processing Unit）的Dispatch架构是其内部处理单元组织和工作流程的重要组成部分，它负责将任务分解为更小的工作单元（通常是线程束或Warp），并将其调度到流多处理器（Streaming Multiprocessors, SM）上进行并行执行。- Warp Scheduler依据一定的策略选择待执行的Warp，这些Warp可能来自同一个kernel的不同线程块，也可能来自不同的kernel。每个SM都有自己的寄存器文件，供其上的线程使用。

2024-04-17 16:06:27 419

原创 L1 CACHE

**缓存控制器**：负责缓存的管理和维护，包括缓存替换策略（如 LRU、Pseudo-LRU 等）、缓存一致性（在多核处理器中，通过 MESI、MOESI 等协议保证数据一致性）、预取机制（预测性地将未来可能需要的数据提前加载到缓存中）等。- **应用程序优化**：程序员应遵循缓存友好的编程实践，如利用数据 locality（工作集大小应适合 L1D Cache 容量）、避免 false sharing（多个线程同时修改共享缓存行的不同部分导致无效的缓存同步）等。L1D Cache 的带宽。

2024-04-17 15:18:12 271

原创 bank冲突

1. **数据访问模式的设计**：通过精心设计数据的访问模式，确保来自同一线程块的线程访问共享内存时，尽可能分散到不同的banks上。5. **广播机制**：当多个线程访问共享内存中的同一地址时，GPU可以采用广播机制，只读取一次数据然后复制给所有请求的线程，而不是每个线程都进行单独的读取操作。7. **优化线程块大小**：选择合适的线程块大小，可以减少线程块内的数据访问冲突。8. **避免连续访问**：在数据加载和访问时，避免连续的线程访问连续的内存地址，这样可以减少bank冲突的发生。

2024-04-17 13:27:49 280

原创 Noc芯片

4. **拓扑结构（Topology）**：定义了路由器和通道的布局，常见的有网格（Mesh）、环形（Ring）和torus等。1. **模块化和可扩展性**：NoC允许系统设计者通过增加更多的节点和连接来扩展系统，而不需要重新设计整个系统架构。2. **提高通信效率**：NoC通过提供多个通信路径和避免共享资源的竞争，提高了数据传输的效率和吞吐量。5. **容错性**：NoC的设计可以考虑容错机制，提高系统的可靠性和稳定性。- **性能**：NoC通过并行通信和避免资源竞争，提高了整体系统的性能。

2024-04-17 12:36:54 322

原创 cross bar

**Intra-Cluster Coalescing**：通过将同一cluster内的SMs的请求合并，减少对crossbar的压力。- **Distributed-Block Scheduling**：通过分布式的块调度策略，优化SMs对内存的访问模式，减少冲突和拥塞。- **缓存策略优化**：通过优化缓存的使用，比如使用SC-Table来解决缓存争用和污染问题，减少因缓存未命中导致的NoC拥塞。- **带宽管理**：通过带宽分配和流量控制策略，确保关键任务有足够的带宽，同时避免不必要的带宽浪费。

2024-04-17 12:31:02 325

原创 tensorcore

然而，这并不会充分利用warp的SIMD特性，因为warp中的线程数量（32）并不能整除矩阵中的元素数量（128），这会导致某些warp中的线程闲置。如果矩阵的大小恰好是warp大小的整数倍（例如，32x32的矩阵），那么每个warp可以自然地处理矩阵的一个部分，而无需额外的操作或考虑。合理的线程排布可以提高访存效率，因为当多个线程访问相邻或相同的数据时，如果这些线程能够在物理上接近的处理器核心上运行，那么它们共享缓存和数据局部性的机会就会增加，从而减少远程内存访问的延迟和开销。

2024-03-27 18:15:27 993

原创 llvm-project

此外，不同的编译器和目标体系结构可能采用不同的寄存器分配策略和优化技术。通过进行过程间的数据流分析，编译器可以确定哪些变量在过程之间是活跃的，哪些变量是死亡的。这些头文件定义了 LLVM 的各种 API 和数据结构，以便其他代码可以使用 LLVM 的功能。因此，当你使用 LLVM 的某个功能时，你通常需要同时引用相应的头文件和链接相应的库文件。文件夹中的头文件定义了 LLVM 的 API 和数据结构，这些 API 和数据结构在。是几个重要的目录，它们分别存放了 LLVM 的头文件、库文件和运行时环境。

2024-03-27 18:12:49 527

原创 cp.async.bulk

是PTX指令集中用于执行异步批量数据搬运操作的指令系列。这类指令允许GPU在后台进行数据的读取或写入，从而尽可能减少等待内存操作完成所带来的延迟，提高执行效率。异步操作的好处是，发起指令的线程可以继续执行其他任务，而非阻塞等待内存操作完成。指令族在GPU编程中扮演着至关重要的角色，它有助于实现数据的高效搬运和处理，优化多线程并行计算的性能。

2024-03-27 11:49:33 308

原创 PTX 8.4 instr

详细说明了在PTX中如何使用。

2024-03-26 18:39:49 693

原创 code emit-mcinstr

对象是通过MCDisassembler从二进制的机器码中解码得到的，并且会包含诸如操作码、操作数、以及可能的隐式信息等细节，这些信息随后可以被MCInstPrinter用来生成汇编语言的表示。它不是一个文本格式，而是一个在LLVM内部使用的编程接口（API）。这段代码只是为了展示如何使用LLVM的API来初始化必要的目标信息并尝试构建一个。通常是在LLVM的内部处理过程中动态创建的，并且与特定的目标架构紧密相关。的内容，我们通常会通过MCInstPrinter这样的类来将。的指令例子是比较困难的，因为。

2024-03-25 19:22:33 284

原创 CallingConv

在LLVM中，函数调用约定（Calling Convention，简称CC）定义了函数如何接收参数以及返回值，包括参数和返回值在内存和寄存器中的布局、由谁负责清理栈等。：根据参数或返回值的类型来选择不同的调用约定。例如，某些类型可能更适合通过寄存器传递，而其他类型可能更适合在栈上分配。例如，一个高级别的函数可能决定其内部调用的低级别函数应该使用哪种调用约定。在LLVM的实际使用中，函数调用约定通常是通过特定的枚举值来定义的，如。通常只在特定的上下文和平台中使用，例如在性能要求较高的内部函数调用中。

2024-03-25 19:19:05 145

原创 alloc register

然而，编译器可能在寄存器分配之后执行其他类型的优化传递，这些传递可能涉及对机器代码的微调、指令重排（以利用处理器的指令级并行性）、窥孔优化（peephole optimization，针对特定指令序列的局部优化）等。尽管这通常不是重写传递的一部分，但值得一提的是，寄存器分配阶段会尝试优化虚拟寄存器到物理寄存器的映射，以减少溢出和不必要的内存访问。虚拟寄存器的数量通常比目标机器上的物理寄存器要多得多。在寄存器分配阶段，编译器会尝试将这些虚拟寄存器映射到有限的物理寄存器上，以便生成可在目标机器上执行的代码。

2024-03-25 17:21:49 500

原创 instruction schedule

在这个例子中，指令I2依赖于指令I1的结果（R1的值），而指令I4依赖于指令I2和I3的结果（R2和R3的值）。例如，当数据需要在不同的函数或程序段之间传递时，或者当编译器进行某些优化操作（如循环展开或指令重排）时，可能需要额外的移动指令来确保数据的正确性和一致性。然后，在寄存器分配阶段，编译器会将这些优化后的指令映射到有限的物理寄存器上，以生成最终的机器代码。虽然这些移动指令可能会增加程序的指令数量和执行时间，但通过合理的优化和寄存器分配策略，可以最大限度地减少这种开销，从而提高程序的性能。

2024-03-25 15:06:51 914

原创 selection instruction

但是，对于加法操作来说，通常不需要引入平台特定的ISD节点，因为加法是大多数平台都直接支持的通用操作。在实际的代码生成过程中，加法操作可能不会经历从通用ISD节点到平台相关ISD节点的转换，而是直接从通用ISD节点转换为平台指令。初始的DAG可能主要由一般LLVM ISD节点构成，随着代码生成过程的进行，这些节点会逐渐被降低为平台相关ISD节点和平台指令。方法是指令选择过程的一部分，它将高级别的中间表示（如SelectionDAG中的ISD节点或平台相关的ISD节点）转换为更低级别的机器指令。

2024-03-25 13:52:47 477

原创开源GPGPU

RTL模型：RTL（寄存器传输级）模型是电子系统设计中的一个较低层次的模型，它描述了数据在寄存器之间的流动和逻辑操作。C模型：C模型是一个较高层次的模型，使用C语言或类似的高级语言来描述系统的行为。在GPGPU中，线程的调度和管理是通过GPU的硬件和软件架构来实现的。开发人员可以通过编程模型提供的API来指定线程块的大小、线程的索引等信息，以便在GPU上正确地调度和管理线程。: 如前所述，Mesa不是一个完整的GPU架构，但它是一个开源的图形驱动库，提供了OpenGL的实现，并支持多种GPU。

2024-03-22 12:01:33 806

原创 llvm后端

SelectionDAGBuilder是LLVM（Low Level Virtual Machine）编译器中的一个重要组件，它负责将LLVM中间表示（Intermediate Representation，IR）转换为SelectionDAG（选择有向无环图）的形式。之前我们分析ir都是抽象出value和use，并且每一条ir都是从clang的ast遍历而来（use关系在ast中如何表示？

2024-03-21 18:05:48 1097

原创 CUDA内置变量—线程索引，内存位置，后端lower，硬件特性

此外，对于某些特定的内置函数，特别是那些与硬件特性或特定优化相关的函数，你可能还需要对目标平台的架构和指令集有深入的了解。作为C/C++编程中作为内置变量，不是直接对应硬件指令，而是CUDA编程模型中的抽象概念，用于标识线程在网格（grid）和块（block）中的位置，帮助开发者组织和标识线程。值，用于计算线程的全局索引，从而访问全局内存中的数据，或者用于执行特定的任务。当你调用内核函数时，CUDA运行时系统会根据你提供的网格和线程块尺寸，自动计算需要的线程总数，并创建相应的线程。

2024-03-21 10:00:40 1166

原创 llvm pass系列

在LLVM中，一致性分析框架并不是一个特定的、官方命名的组件，但我们可以理解为在LLVM的架构和工具集中，存在一系列用于确保代码转换和优化过程中一致性的机制和方法。可能除了ir和pass之外，对编译器语义不变这个核心的理论，也值得参考学习，面试的时候最多也就问你知道哪些pass，几乎不会问这些问题，但是这是编译器原理，是编译器的目标约束。它们基于分析pass提供的信息，对程序进行实际的转换和优化操作。LLVM的pass架构使得相同的优化和分析可以在不同的目标架构和操作系统上运行，提高了代码的跨平台兼容性。

2024-03-21 09:39:58 435

原创 ptx指令，抽象指令与架构

了解PTX指令集可以为你在编译器学习旅程中提供一个更深层次的视角，让你能够更好地理解、优化和利用GPU计算的潜力。了解现有的PTX指令集可以让你为未来的变化和扩展做好准备，使你能够更好地利用新架构中的新功能和优化。了解PTX可以帮助你更深入地理解并行计算是如何在现代GPU上实现的，包括线程的组织、同步和内存访问模式等。了解PTX可以让你更好地理解底层硬件是如何工作的，包括内存层次结构、计算单元的组织、指令流水线和调度策略等。谓词操作数通常表示为以。: 其他同步操作的指令，可能用于确保内存操作的顺序一致性。

2024-03-20 15:58:19 1350

原创 llvm IR指令

在实际的代码生成阶段，Phi指令会被转换成具体的机器指令序列，这些指令会根据控制流条件来选择正确的值。在LLVM IR中，Phi（φ）指令是SSA（静态单赋值）形式的关键组成部分。ir之间，因为是ssa形式，每一个变量的创建和被变成operand使用都有一个uses结构（数组或者链表），在定义时，该变量就继承并初始化了uses，这对以编译器掌握每个变量的生存周期是重要的基础。LLVM IR中有很多种不同的Instruction，比如load，store，add，call等等，每种指令用于实现特定的功能。

2024-03-20 14:09:25 746

原创编译器的优化选项和简单解析

编译器的优化级别O0、O1、O2和O3在代码编译时提供了不同的优化程度，它们的主要区别在于编译时间、目标文件大小和执行效率之间的平衡。除了以上四个优化级别外，还有一些扩展的优化级别，如Os（优化代码尺寸）。这个级别旨在优化代码的尺寸，尽量减小可执行文件的大小。它对于磁盘空间紧张或CPU缓存较小的机器非常有用，但也可能产生些许问题。因此，在软件树中的大部分情况下并不推荐使用Os优化级别。在计算机科学中，控制流图（Control Flow Graph，CFG）是一个表示程序所有可能执行路径的抽象数据结构。

2024-03-18 19:49:29 2262

原创 NPU、GPU的数据路径

优化NPU的数据路径是提高神经网络算法性能的关键步骤之一。这里的描述可以直观的理解tvm和tvm向上的推理引擎，云计算框架都可叫编译器，向下的llvm，甚至微架构设计，都可以叫编译器，只是处理的问任务，问题，计算，微计算的粒度层级不一样。之前面试的时候，遇到：请描述你所知架构的数据路径，挺懵，不知到是从哪些方面说。在实际的NPU设计中，数据路径可能更加复杂，包括多级存储层次结构、多个并行处理单元、复杂的数据依赖关系和数据重用策略等。一般面试问这么大的问题，而对问题本身又没有约束的，面试官可能真是个官吧。

2024-03-18 19:49:18 400

原创 @llvm.amdgcn.workitem.id.x()引发的一些前后端的调研

（这不又是intrinsic函数嘛，但是它又有自己的指令集，又有很多intrinsic函数，比如矩阵乘法，如果是这样的话，它和llvm ir确实也是一样的形式，都有指令和函数。在OpenCL的上下文中，运行时编译可以看作是一种动态编译的形式，因为它允许在程序运行时动态地生成和编译内核代码。虽然LLVM IR是LLVM项目的一种通用中间表示形式，用于在各种编译器前端和后端之间进行转换，但在处理CUDA代码时，Clang选择使用NVVM IR而不是LLVM IR，因为NVVM IR更适合GPU代码的特殊需求。

2024-03-13 17:33:46 846

u012276729的博客