CUDA-interview-note

CUDA-interview-note

Charles Lv7
  2024-04-15 18:10:17 2024-04-15 18:10:17 Created   2024-04-15 18:28:09 2024-04-15 18:28:09 Updated      19.3k Words  67 Mins  

转自：国内大厂GPU CUDA高频面试问题汇总（含部分答案） - 知乎 (zhihu.com)

整体情况简介

面试中的问题基本上分成以下几类：

  1. 基础的八股文：C/C++，OS，计算机体系结构等。
  2. 高性能计算基础知识：这一部分是面试的重点，本文章以CUDA为重点。
  3. 各种AI框架知识：本文章以推理方向为主。
  4. AI基础知识：对于常见的机器学习算法，以及CV & NLP & 推荐模型有一定了解，了解计算流程以及模型结构即可，重点为了能分析出计算瓶颈在哪里，找出可能优化的方向。本部分略
  5. 算法题： 手写CUDA kernel和leetcode的比例大约为3:1。手写CUDA kernel的时候一般会结合第2部分一起问，一步一步要求你优化，每一步优化的具体原理，涉及到什么硬件知识等。

高性能计算基础

  1. CUDA的线程组织结构
  2. CUDA的存储体系结构，每一种存储的优缺点，该如何合理使用。
  3. GPU每一代的新特性有了解过吗？应该从哪里去了解详细信息？
  4. CUDA stream的概念，为什么要使用多个stream？
  5. GPU和CPU分别适合执行哪些程序？结合它们的硬件架构解释一下为什么它们有各自的优势。
  6. 说明一下神经网络加速器与CPU、GPU的区别，他们各自有何优势？
  7. 半精度浮点数FP16各个部分的具体位数，为什么要有半精度浮点数？
  8. TensorCore的加速原理
  9. MPI，OpenMP以及CUDA各自适用的加速场景。
  10. RDMA相关问题。
  11. 平时如何进行kernel的优化，会用到哪些工具？
  12. CPU上哪些并行优化方法？
  13. ARM相关的库有了解过吗？
  14. PTX有了解过吗？
  15. roofline模型有什么用？如何确定最优的BLOCK_SIZE。
  16. GPU资源调度有哪些方法？
  17. 稀疏矩阵的存储格式有哪些？稀疏矩阵的应用场景？稀疏矩阵计算与稠密矩阵计算有何不同？
  18. 如何计算CPU指令的吞吐量和时延?

AI 框架知识

这一部分会涉及一些AI框架(训练&推理&编译器)相关的问题，并且会重点根据简历上的项目经历去做一些发散性的提问。

  1. MLIR有了解过吗？ONNX有了解过吗？
  2. TVM的整体结构，如何用TVM进行开发？
  3. 为什么要进行推理优化？直接用tensorflow或者pytorch的推理接口不行吗？
  4. 模型推理优化的常用方法有哪些？
  5. 有研究过某一个框架的具体源码吗？
  6. TensorRT如何进行自定义算子开发？
  7. TensorRT对模型实现了哪些推理优化？常量折叠，算子融合，量化…
  8. 算子融合为什么能加速推理，优化了哪一部分？TensorRT用到了哪些算子融合？算子融合在推理框架中是如何实现的？
  9. 模型量化的加速原理，模型量化带来的精度损失如何解决？
  10. ONNX Runtime支持在多种硬件上进行推理，说明具体的实现机制。
  11. 总结一下TensorRT，ONNX Runtime等推理框架的组成架构，如果我们公司自己要为硬件开发一套推理框架，应该重点关注哪些部分？
  12. 各种推理框架都有何优劣势？它们的性能怎么样？
  13. 分布式训练中有哪些并行模式？每种模式需要做什么，有什么优缺点？
  14. 分布式训练中我们重点需要处理的问题有哪些？目前已有哪些解决方案
  15. MPI如何应用于AI框架中？
  16. 模型在移动端进行推理优化的框架有了解过吗？移动端和在服务器的推理优化思路有何不同？移动端能用到的加速指令有了解过吗？
  17. 移动端有哪些加速方法？
  18. 为什么要将模型一部分推理优化放在移动端，全部放在服务器上不可以吗？
  19. 自动驾驶上的推理框架有了解过吗？我们重点需要关注的指标有哪些？
  20. 反向传播的原理，具体实现的源码有了解过吗？
  21. 你了解哪些推理模型的调度方法？
  22. 推荐模型的结构有了解过吗？要部署一个大的推荐模型，应该如何将各个部分放在哪种硬件上部署？
  23. 计算图切分有了解过吗？如何应用于大模型推理？
  24. TensorFlow和Pytorch都用过吗？它们设计思路有何不同？有何优劣？如何添加自定义算子？

算法题

手写CUDA kernel几乎每场面试都会考，面试官会以写出来的第一个版本为准，一步步问继续优化的方法，在这个期间会结合高性能计算的基础知识来考察，从这个过程中能了解到对体系结构以及优化方法的了解程度。leetcode不一定有，但是遇上了基本上都是hard。两类算法题都要准备。

下面是常见的一些问题：

  1. 矩阵乘:
  2. 矩阵转置: 访存密集型算子的处理
  3. 一维reduce-sum：重点是如何处理bank confict
  4. 二维reduce-sum
  5. 卷积
  6. 将单stream改成多stream

以矩阵乘法为例说明一下一个典型的面试流程，下面以A表示面试官，B表示面试者。

A：写一个矩阵乘法吧，并将main函数中具体调用给写清

B: （写了一个最naive版本的矩阵乘）

A: 目前这个程序有什么问题，能进一步优化吗？

B : 目前访存性能比较低，可以采用矩阵分块并且使用上shared memory优化，并解释一下这样做的原理。

A：可以具体计算一下优化前后的计算访存比，来具体说明这一部分提升了多少。并写一下优化后的程序。

B: 通过计算优化了…

上述对话会重复几轮，在后面几轮可能面试官不会再要求将每一版程序都写出来了，重点在于讨论优化思路，并且在讨论的过程中发散地问一点CUDA的知识考察理解的深度。

一些比较零碎的问题

  1. 卷积的三种加速计算方式，im2col+GEMM & Winograd & FFT，各自有何优缺点，cuDNN中有哪些实现？
  2. 数字信号的采样定理、熵 & 交叉熵 的含义 & 计算公式
  3. 还记得KKT条件吗？写程序求解一个非线性方程，并说明具体用到的优化方法。
  4. 脑洞问题：如何从编码的角度进行模型压缩？
  5. 如何将你研究生阶段的成果应用到我们的产品中？
  6. 给了一个TF 模型的profile，找出里面的bottle neck，提出如何改进这个模型的性能的方法。
  7. MIPS流水线有几级？分别是哪些组成部分？
  8. 说一下transformer的具体结构，如何加速transformer进行推理？
  9. attention的计算公式，写一下tf里面对应的代码
  10. 马尔科夫链简单知识
  11. 一道较难的概率题

推荐参考资料

  1. 《通用图形处理器设计:GPGPU编程模型与架构原理》：CUDA、GPU体系结构、PTX、TensorCore等GPU知识大杂烩，CUDA相关面试问题标答。对于GPU的硬件体系结构有较深入的介绍，虽然比较难懂，但是这一部读完后会对CUDA编程模型以及为什么要采用一些特定的优化方法有更深入的理解。
  2. 官方文档《CUDA Programming Guide》 & 《CUDA Best Practice Guide》: 不解释，必读。
  3. 《大规模并行处理器程序设计》：入门最佳，没有之一。其中第二部分对于CUDA中常见的计算Pattern做了分析，几乎可以应付所有的面试中的kernel编程，至少能答出80%，至于更深入地优化方法需要再花时间去研究。
  4. 《机器学习系统：设计和实现》：介绍了ML Sys这一领域的所有方面的基础知识，可以从一个整体的层面对机器学习系统的组成部分、每个部分的重点技术有较好的把握。这本书的框架主要以MindSpore为例，所以在整体读完后，需要结合自己比较熟悉的框架进一步仔细理解。该书有在线版本机器学习系统：设计和实现 - 机器学习系统：设计和实现 1.0.0 documentation
  5. 《深度学习进阶：自然语言处理》：只用numpy实现NLP模型，可以作为阅读深度学习框架源码的first course，会对AI模型中的底层实现细节有很好的理解。
  6. 《分布式机器学习：理论、算法与实践》：可以对分布式训练有大致的了解
  7. 《AI编译器开发指南》：深度学习编译器相关的介绍，重点在TVM。

建议： 1 ~ 4必读，这是所有领域的基础知识，5 ~ 7需要根据个人的研究兴趣和方向有选择性地深入阅读。

高性能计算基础

1. 你可以简单说下CUDA的线程组织结构吗？

相对于GPU硬件结构来说，CUDA是一个并行计算的编程模型，它是基于GPU的体系结构设计的。为了高效地利用GPU的并行计算能力，简化并行编程，提高程序的性能和可扩展性，CUDA提供了一种抽象的层次模型，即CUDA的线程组织结构。

CUDA的线程组织结构包括网格（grid）、线程块（block）和线程（thread）。

首先，GPU核心程序 kernel 在 device上 执行时启动很多 Thread，而一个 kernel 所启动的所有线程称为一个网格 grid。

其次，同一个 grid 上的线程共享相同的全局内存空间，而 grid 又可以分为很多线程块 block，线程块是向GPU进行调度的最小单位，GPU同时支持多个线程块的执行，达到上限后，只有旧的线程块内的线程全部执行完成后，新的线程块才会被调度入GPU。

最后，一个 block 里面包含很多 Thread，每个线程Thread都是独立执行的，并且可以访问全局内存和共享内存，线程之间可以通过同步原语进行同步操作。

在进行CUDA编写gpu kernel函数时，需要使用__global__修饰，返回值必须为void 。在核函数的调用时需要使用三括号的方式来指明核函数中的线程数目以及排列情况的，如<<<1, 1>>>

三括号中的第一个数字可以看作线程块的个数，第二个数字可以看作每个线程块中的线程数，三括号中的两个数字分别就是网格大小和线程块大小，即 <<<网格大小, 线程块大小>>> 。而核函数的总线程数即为网格大小乘以线程块大小。

值得注意的是从开普勒架构开始，最大允许的线程块大小是1024。

img

此外，grid和block在CUDA的编程模型中也支持多维组织。

1
2
3
dim3 grid(2, 2);
dim3 block(4, 2, 2);
kernel_fun<<< grid, block >>>(prams...);

如上所示，grid 和 block 都是定义为 dim3 类型的变量。

CUDA的这种 <<<grid,block>>> 其实就是一个多级索引的方法，第一级索引是 (grid.xIdx, grid.yIdy)，对应上图选中的 Block 就是(1, 1)，通过它我们就能找到了这个线程块的位置，然后我们启动二级索引 (block.xIdx, block.yIdx, block.zIdx) 来定位到指定的 Thread。这就是CUDA的线程组织结构。

img

img

从硬件上来看，CUDA编程模型中的网格、线程块和线程大致与硬件结构中的GPU、SM（流式多处理器）和SP（流式处理器）是一一对应的。

img

从上图可以看出，一个Grid可以包括多个SM，也可以访问Global Memory和Constant Memory；

一个Block只能在一个SM中，且一个SM包含多个Block，Block可以访问Shared Memory；

一个Block中有多个Thread，而一个Thread只能访问Registers或local Memory。

此外，一个线程块Block还可以细分成多个线程束，一个线程束（也就是一束线程）是一个线程块里面相邻的warpSize个线程，目前warpSize 都为32。

一般来说，希望线程块的大小是warpSize的整数倍，否则系统会自动为剩下的n个线程补齐32-n个线程，形成一个完整的线程束，而这32-n个线程并不会被核函数调用，从而闲置。

2. 可以简单谈谈CUDA的存储体系结构，每一种存储的优缺点，该如何合理使用。

CUDA的存储体系结构包括全局内存（Global Memory）、共享内存（Shared Memory）、常量内存（Constant Memory）/纹理内存（Texture Memory）和本地内存（Local Memory）。

img

  1. 全局内存（Global Memory）

这是GPU中最大的内存，即我们常说的HBM内存，可以被所有块上的所有线程访问，当我们在GPU中初始化一个值而不指定其存储位置时，它会自动存储在全局内存中。然而，访问全局内存通常比其他内存类型慢，因此需要进行优化以避免性能下降，可以通过合并内存访问和使用共享内存来优化性能。

2. 共享内存（Shared Memory）

同一个Block内的线程可以通过共享内存共享数据。相比访问全局内存至少快个10倍，但共享内存的容量有限（通常为几十KB），无法被其他线程块访问。由于共享内存和缓存内存提供快速的访问速度，因此我们经常在计算过程中使用它们来存储数据。典型的方法是首先将所有数据从 CPU 复制到 GPU 并将其存储在全局内存中。然后，我们将数据分解成更小的部分（块）并将它们推送到共享内存中进行计算。计算完成后，结果将被推回全局内存。

3. 纹理内存和常量内存（Texture and Constant Memory）：

这些是GPU 中的特殊内存类型，针对访问特定数据类型（例如纹理或常量值）进行了优化。所有块中的所有线程都可以访问这些内存类型。

例如，常量内存专门只能用于存储只读数据，纹理内存只能用于存储二维图像数据，这两种内存类型的访问速度都相当快，可以与共享内存相媲美。

因此，使用纹理内存和常量内存的目的是优化数据访问并减少共享内存的计算负载。我们可以将一部分数据分配给纹理内存和常量内存，而不是将所有数据推送到共享内存中。这种分配策略通过利用纹理内存和常量内存的优化访问功能来帮助增强内存性能。

4. 本地内存（Local Memory）

每个线程都可以使用自己的本地内存，可以在其中存储临时变量。它具有最小的范围并且专用于每个单独的线程。

3. 你了解CUDA stream吗？为什么要使用多个stream？

stream相当于是GPU上的任务队列，用官方的话叫做一条命令流水线，它允许多个CUDA操作在不同的stream中并行执行，从而提高GPU的利用率和性能。

每个kernel调用或大多数CUDA API都可以指定关联到某一个stream，同一个stream的任务是严格保证顺序的，上一个命令执行完成才会执行下一个命令。

不同stream的命令不保证任何执行顺序，部分优化技巧需要用到多个stream才能实现。如在执行kernel的同时进行数据拷贝，需要一个stream执行kernel，另一个stream进行数据拷贝，或者针对大数据集进行切分，然后可以采用多个stream执行并行加速拷贝。

此外，多个stream还能够方便地划分和管理不同的任务，提高应用程序的灵活性和可扩展性。

img

4. GPU和CPU分别适合执行哪些程序？结合它们的硬件架构解释一下为什么它们有各自的优势。

img

CPU 将大量芯片面积专门用于可减少指令延迟的功能，例如大缓存、更少的 ALU 和更多的控制单元。

相比之下，GPU 使用大量 ALU 来最大化其计算能力和吞吐量，它们使用非常少量的芯片区域作为缓存和控制单元，使得其具有很高的延迟。

所以对于他们的设计目标来说，CPU 被设计为顺序执行指令，为了提高顺序执行性能，多年来 CPU 设计中引入了许多功能，包括指令流水线、乱序执行、推测执行和多级缓存等。CPU的重点是减少指令执行延迟，以便 CPU 能够尽快执行指令序列。

GPU 专为大规模并行性和高吞吐量而设计，但代价是较高的指令延迟。这一设计方向受到了它们在视频游戏、图形、数值计算和现在深度学习中的使用的影响。所有这些应用程序都需要以非常快的速度执行大量线性代数和数值计算，因此人们对提高这些设备的吞吐量投入了大量注意力。

此外，GPU拥有高带宽的内存和专门用于并行计算的指令集，能够更高效地处理大规模数据的并行计算任务。

5. 说明一下神经网络加速器与CPU、GPU的区别，他们各自有何优势？

在CPU中70%晶体管用来构建Cache，还有一部分控制单元，计算单元少，所以说CPU的核心擅长完成多重复杂任务，重在逻辑，重在串行程序；

GPU的计算模型是单指令、多数据SIMT处理，晶体管大部分构建计算单元，运算复杂度低，适合大规模并行计算GPU的核心擅长完成具有简单的控制逻辑的任务，重在计算，重在并行，适合深度学习，图像处理，大数据领域。但GPU无法单独工作，必须由CPU进行控制调用才能工作。

NPU，即神经网络加速器，在电路层模拟神经元，相比于CPU和GPU，NPU通过突出权重实现存储和计算一体化，从而提高运行效率，一条指令完成一组神经元的处理，提高运行效率。NPU是模仿生物神经网络而构建的，CPU、GPU处理器需要用数千条指令完成的神经元处理，NPU只要一条或几条就能完成，因此在深度学习的处理效率方面优势明显，但它需要专用定制化实现，而且不能随意扩展，并不通用。

img

6. 半精度浮点数FP16各个部分的具体位数，为什么要有半精度浮点数？

半精度浮点数（FP16）是一种二进制浮点数格式，其具体位数如下：

  • 1位符号位：用于表示正负号。
  • 5位指数位：用于表示指数部分，取值范围为-14至15（实际上是通过偏移值来表示，即真实的指数值减去15）。
  • 10位尾数位：用于表示尾数部分。
  • float16 最大范围是 [-65504 - 66504]
  • float16 能表示的精度范围是 2^−24 ，超过这个数值的数字会被直接置0；

float16和float相比恰里，总结下来就是两个原因：内存占用更少，计算更快。

  • 内存占用更少：这个是显然可见的，通用的模型 fp16 占用的内存只需原来的一半。memory-bandwidth 减半所带来的好处：

    • 模型占用的内存更小，训练的时候可以用更大的batchsize。
    • 模型训练时，通信量（特别是多卡，或者多机多卡）大幅减少，大幅减少等待时间，加快数据的流通。

  • 计算更快：

    • 目前的不少GPU硬件都有针对 fp16 的计算进行优化，例如TensorCore，半精度的计算吞吐量可以是单精度的 2-8 倍；

那既然fp16像上面说的那么好，那么是否全部都使用 fp16 即可了呢？

当然不是，全用fp16 主要存在两个问题：1. 数据溢出问题；2. 舍入误差。

7. 可以谈下TensorCore的加速原理吗？

首先，当谈到Tensor Core的计算速度时，与CUDA Core相比，它能够在一个时钟周期内执行多个操作。Tensor Core可以同时对两个4×4的FP16张量进行矩阵乘积计算，并将结果累加到另一个4×4的张量上（即D = A * B + C）。

其次，TensorCore的加速也基于混合精度矩阵乘法。混合精度并不是网络层面既有 FP16 又有 FP32，它指的是在底层硬件算子层面，使用半精度（FP16）作为输入和输出，使用全精度（FP32）进行中间结果计算从而不损失过多精度的技术。

通过硬件上的特殊设计，Tensor Core 理论上可以实现 8 倍于 FP32 Core 的计算吞吐量（Volta 和 Turing 架构），并且没有明显的占用面积和功耗增加。

img

既然Tensor Core这么好，为什么Nv不生产一个全部是Tensor Core的深度学习专用GPU架构？

虽然说Tensor Core是专门为加速深度学习和 AI 工作负载（例如矩阵运算）而设计，但目前深度学习也不能离开Cuda Core。

CUDA Core 针对各种并行计算任务进行了优化，更适合于通用并行计算任务。

首先，深度学习任务其不仅仅是矩阵运算，还有很多的并行计算，这就看瓶颈在哪里了。如果瓶颈时在并行计算，那么这种类型的深度学习任务可能更适合CUDA Core。

其次，Tensor Core的使用是有些限制的，对于GEMM计算其效果很好，其次其输入和输出数据类型需要是半精度或单精度，矩阵的维度最好是 8 的倍数。

当然Tensor Core也可用于通用并行计算任务，但它们可能不如 CUDA Core 高效。

8. MPI，OpenMP以及CUDA各自适用的加速场景。

  1. MPI

MPI是一个跨语言的通讯协议，支持高效方便的点对点、广播和组播。从概念上讲，MPI应该属于OSI参考模型的第五层或者更高，他的实现可能通过传输层的sockets和TCP覆盖大部分的层。

MPI是基于消息传递的并行编程的，用户程序利用这些接口进行进程之间的数据移动、聚集、规约和同步。MPI标准规定了这些接口的调用规范和语义，不同的实现（例如mpich或者openmpi）可能采用不同的优化策略。

其中，点对点通信指的是两个进程之间的通信，可用于控制同步或者数据传输，例如MPI_Send和MPI_Recv。

集合通信包括了一对多、多对一和多对多的通信方式，常用于一组进程之间的数据交换，例如AlltoAll，Allreduce等。

它适用于需要在大规模分布式系统上进行高性能计算的场景，如集群计算等。

img

既然谈到了集合通信，可以谈一下Ring-Allreduce吗，为什么深度学习需要Ring-Allreduce？

假设有5块GPU，每一块GPU拥有完整的模型参数可以进行forward pass和backward pass，总共的训练数据大小为K，我们需要根据训练数据计算出所需要的梯度进行一次迭代。

考虑一个简单的同步通信策略。首先，每张GPU拥有同样的初始参数，我们将大小为K的训练数据分为N块，也就是5块，分给每张GPU。每个GPU基于自己那一部分的数据，计算得到本地的local gradients，然后N-1块（4块）GPU将计算所得的local gradients 发送给GPU 0，让GPU 0对所有的local gradients进行reduce(汇聚操作)得到全局的梯度，然后再将该全局梯度返回给每块GPU进行back propagation来更新每个GPU上的模型参数。

那么我们可以计算下整个通信过程的communication cost。首先要记住现实中，network的 bandwidth是有限的，假设每张GPU需要发送给GPU 0的通信数据大小是1GB，我们的network bandwidth是1GB每秒（GPU 0最多每秒接受1GB的数据），那么我们需要4秒才可以将数据全部发送到GPU 0上，然后计算出全局的平均梯度。我们可以看到其通信成本是 C * N，由于受GPU 0的network bandwidth的影响，通信成本随着设备数的增加，而线性增长。

img

相比之下，Ring Allreduce的通信成本恒定，和设备数量无关，完全由系统中GPU之间最慢的连接决定。

我们将所有设备安排在一个逻辑环中，每个GPU应该有一个左邻和一个右邻，设备只会从它的右邻居发送数据，并从它的左邻居接收数据，整个计算过程通过Scatter reduce和Allgather两个通信原语完成。

\2. OpenMP

OpenMP 是基于共享内存模式的一种并行编程模型。

OpenMP 是以线程为基础的，其执行模式采用fork-join的方式，其中fork创建新线程或者唤醒已有的线程，join将多个线程合并。

在程序执行的时候，只有主线程在运行，当遇到需要并行计算的区域，会派生出线程来并行执行, 在并行执行的时候, 主线程和派生线程共同工作, 在并行代码结束后, 派生线程退出或者挂起，不再工作，控制流程回到单独的线程中。

img

OpenMP适用于单台计算机上的多核并行计算。通过在代码中插入指令，开发者可以指示并行执行，并将任务分配给多个处理器核心。OpenMP适用于需要在单个计算节点上进行并行计算的场景，如多核处理器、多线程编程等。

\3. CUDA

CUDA（Compute Unified Device Architecture）是一种用于GPU加速计算的并行计算平台和编程模型。CUDA适用于利用GPU进行并行计算的场景。

通过编写CUDA C/C++代码，开发者可以将计算任务分配给GPU上的成百上千个并行计算单元（CUDA核心），以实现高效的并行计算。CUDA适用于需要大规模并行计算的科学计算、机器学习、深度学习等领域。

9. 可以说下DMA和RDMA是什么吗？以及有哪些硬件上的实现。

首先，DMA(直接内存访问)是一种能力，允许在计算机主板上的设备直接把数据发送到内存中去，数据搬运不需要CPU的参与。

img

传统内存访问需要通过CPU进行数据copy来移动数据，通过CPU将内存中的Buffer1移动到Buffer2中。DMA模式：可以同DMA Engine之间通过硬件将数据从Buffer1移动到Buffer2,而不需要操作系统CPU的参与，大大降低了CPU Copy的开销。

其次，RDMA其实从名字上就可以看出，其多了一个R，即Remote。指的是在两个或者多个计算机进行通讯的时候使用DMA， 从一个主机的内存直接访问另一个主机的内存。

img

RDMA是一种新的直接内存访问技术，RDMA让计算机可以直接存取其他计算机的内存，而不需要经过处理器的处理。RDMA将数据从一个系统快速移动到远程系统的内存中，而不对操作系统造成任何影响。

在实现上，RDMA实际上是一种智能网卡与软件架构充分优化的远端内存直接高速访问技术，通过将RDMA协议固化于硬件(即网卡)上，以及支持Zero-copy和Kernel bypass这两种途径来达到其高性能的远程直接数据存取的目标。 使用RDMA的优势如下：

  • 零拷贝(Zero-copy) - 应用程序能够直接执行数据传输，在不涉及到网络软件栈的情况下。数据能够被直接发送到缓冲区或者能够直接从缓冲区里接收，而不需要被复制到网络层。
  • 内核旁路(Kernel bypass) - 应用程序可以直接在用户态执行数据传输，不需要在内核态与用户态之间做上下文切换。
  • 不需要CPU干预(No CPU involvement) - 应用程序可以访问远程主机内存而不消耗远程主机中的任何CPU。远程主机内存能够被读取而不需要远程主机上的进程（或CPU)参与。远程主机的CPU的缓存(cache)不会被访问的内存内容所填充。
  • 消息基于事务(Message based transactions) - 数据被处理为离散消息而不是流，消除了应用程序将流切割为不同消息/事务的需求。
  • 支持分散/聚合条目(Scatter/gather entries support) - RDMA原生态支持分散/聚合。也就是说，读取多个内存缓冲区然后作为一个流发出去或者接收一个流然后写入到多个内存缓冲区里去。

在具体的远程内存读写中，RDMA操作用于读写操作的远程虚拟内存地址包含在RDMA消息中传送，远程应用程序要做的只是在其本地网卡中注册相应的内存缓冲区。远程节点的CPU除在连接建立、注册调用等之外，在整个RDMA数据传输过程中并不提供服务，因此没有带来任何负载。

**最后，RDMA 三种不同的硬件实现。**RDMA作为一种host-offload, host-bypass技术，使低延迟、高带宽的直接的内存到内存的数据通信成为了可能。目前支持RDMA的网络协议有：

  1. InfiniBand(IB): 从一开始就支持RDMA的新一代网络协议。由于这是一种新的网络技术，因此需要支持该技术的网卡和交换机。
  2. RDMA过融合以太网(RoCE): 即RDMA over Ethernet, 允许通过以太网执行RDMA的网络协议。这允许在标准以太网基础架构(交换机)上使用RDMA，只不过网卡必须是支持RoCE的特殊的NIC。
  3. 互联网广域RDMA协议(iWARP): 即RDMA over TCP, 允许通过TCP执行RDMA的网络协议。这允许在标准以太网基础架构(交换机)上使用RDMA，只不过网卡要求是支持iWARP(如果使用CPU offload的话)的NIC。否则，所有iWARP栈都可以在软件中实现，但是失去了大部分的RDMA性能优势。

10. 平时如何进行kernel的优化，会用到哪些工具？

首先，要优化kernel函数需要先了解GPU硬件的构造。其次，需要熟悉常见的profiler工具，主要包括Nsight System和Nsight Compute。

img

在优化的手段和方向上主要关注几个点：

1. 使用异步API

使用异步API如cudaMemcpyAsync可让GPU操作与CPU操作并行，CPU忙完后调用cudaStreamSynchronize，cudaEventWait等操作等待GPU任务完成。

2. 优化内存与显存传输效率

  • 使用Pinned(page-locked) Memory提高传输速度
  • 通过在不同的Stream里同时分别执行kernel调用及数据传输，使数据传输与运算并行。（注意default stream的坑）
  • 尽量将小的数据在GPU端合成大块数据后传输

3. 优化Kernel访存效率

  • 提高Global Memory访存效率
  1. 对Global Memory的访存需要注意合并访存(coalesced )。
  2. warp的访存合并后，起始地址及访存大小对齐到32字节
  3. 尽量避免跨步访存
  4. CUDA 8.0及以上的设备可以通过编程控制L2的访存策略提高L2命中率。
  • 提高Shared Memory的访存效率
  1. shared memory由32个bank组成
  2. 每个bank每时钟周期的带宽为4字节
  3. 连续的4字节单元映射到连续的bank。如0-3字节在bank0，4-7字节在bank1……字节128-131字节在bank0
  4. 若warp中不同的线程访问相同的bank，则会发生bank冲突(bank conflict)，bank冲突时，warp的一条访存指令会被拆分为n条不冲突的访存请求，降低shared memory的有效带宽。所以需要尽量避免bank冲突。
  5. CUDA 11.0以上可以使用async-copy feature

4. 优化线程级并行

在SMSP工作时，某些warp会由于访存依赖、寄存器依赖等原因stall。此时warp scheduler可以选中另一个eligible warp，执行其指令，以隐藏前一个warp的stall，使SMSP中的各个硬件资源尽量保持忙碌。但假如SMSP中所有的warp都不在eligible状态，则硬件只能空转等待某个warp从stall中恢复(如从global中请求的数据终于回来了)。

Occupancy指标用来衡量SM当前activate warp数量与理论上最多支持的activate warp数量的比值。Occupancy数量越高，代表SMSP负责的activate warp越多，当某个warp stall时，有更多的备选warp，有更大的概率可以找到一个eligible warp。极端情况Occupancy为1/8时，SM仅4个warp，每个SMSP 1个warp，当该warp stall时，smsp没有其它warp可以选择，硬件必然空转等待。

影响Occupancy指标的包括以下因素：

  1. Thread Block 线程块的大小。
  2. 每个线程块的Shared Memory使用量
  3. 每个线程使用的Register（寄存器数量）

高的Occupancy不一定代表较高的性能，如某些算法确实需要每线程128寄存器时，保持0.5的Occupancy反而是最优选择。但过低的Occupancy会对性能带来较大的负面影响。

5. 指令级优化

  • 提高计算访存比

GPU执行计算时，需要LDS、LDG等指令先将数据读入寄存器，再进行计算，最后通过STS、STG等指令将数据保存下来。

以矩阵乘法为例，先进行矩阵分块，最终拆解为每个线程计算MxK,KxN的两个小矩阵的乘法：

若两小矩阵为M=2,N=2,K=1，即2x1;1x2,最后得到2x2的矩阵作为结果。则读入4个float需4条指令，计算指令也是4条，计算访存比4/4=1；

若两小矩阵为M=8,N=8,K=1，即8x1;1x8,最后得到8x8的矩阵作为结果。则读入16个float，需读取指令16条，计算指令8x8=64条，计算访存比64/16=4；若使用向量读(float4)每条指令读入4个float，则读取指令仅4条，计算访存比64/4=16

提高计算访存比，可以让GPU的更多时钟周期用于进行计算，相对的进行数据IO占用的时钟周期更少。

  • 提高指令级并行

指令级并行基本原理：

  • 现代不论是CPU还是GPU，指令的执行都是通过流水线进行的，流水线分为多个stage，即一条指令执行完成需要每个stage的工作都执行完成。而一个时钟周期并不是完成一条指令执行的所有时间，而是每一个stage完成当前工作的时间。流水线可以同时执行多条指令的不同阶段。
  • 当后续指令的执行需要依赖前面指令的结果写回寄存器，我们说出现了寄存器依赖。此时后续指令需要等待第前面指令结果写回寄存器才能执行，若后续指令执行时前面指令结果尚未写回寄存器，流水线会失速（stall），此时warp scheduler开始切换到其它eligible warp，若无eligible warp，则SMSP将会空转。
  • 若后续指令不依赖前面指令的结果，则即使前面指令未执行完毕，后续指令也可以开始执行。特别的，即使前序指令是一条耗时几百周期的LDG(全局内存读取)指令或耗时几十周期的LDS（共享内存读取）指令，只要后续一系列指令不依赖读取回来的数据，后续一系列指令可以正常执行而不必等待该LDG/LDS指令执写回寄存器。

通过以下方式，可以提高指令级并行，在线程级并行达不到较好效果的情况下，进一步提高程序性能：

  • 数据预取（Prefetch）：数据1已读取到寄存器，使用该数据1计算前，先将后续数据2的读取指令发射，再执行一系列数据1的处理指令；这样数据1的处理和数据2的读取在流水线上同时执行着。当数据1处理完成，需要处理数据2时，可以确保数据2已经存在于寄存器中，此时类似的将数据3的读取和数据2的处理同步执行起来。
  • 指令重排：在存在寄存器依赖的指令间插入足够的其它指令，使得后续指令执行时，前面计算指令的结果已写回到寄存器。从CUDA C层面有意识地提供一些语句间的并行性，nvcc编译器可以一定程度上自动进行指令重排。若对nvcc重排结果不满意需要自己重排时，官方尚未开放SASS汇编器，目前只存在一些第三方SASS汇编器工具。
  • 提高Register的效率
  1. Register File也存在bank冲突，但在CUDA C层面上没有直接办法进行物理寄存器控制。
  2. 可以通过SASS汇编器，人工进行指令寄存器分配，以尽量消除register bank conflict。
  3. 可以通过SASS汇编器，为寄存器访问添加reuse标记，以尽量消除register bank conflict。

6. 使用TensorCore进一步加速矩阵运算

TensorCore可以用来快速进行D=A*B+C矩阵运算，提供load_matrix_sync， store_matrix_sync， mma_sync 等API。

可参考这篇文章：https://zhuanlan.zhihu.com/p/570795544

11. CPU上哪些并行优化方法？

img

  1. 线程级并行：将任务分解成多个线程，利用多核处理器同时执行这些线程，从而加快任务完成速度。
  2. SIMD指令集：使用单指令多数据的指令集，同时处理多个数据，提高向量运算的效率。
  3. OpenMP和MPI：这是一些常用的并行编程框架，可以帮助开发人员实现并行计算，充分利用多核处理器的性能。
  4. 数据并行：将数据分割成小块，分配给不同的处理器核心并行处理，加快数据处理速度。

img

12. ptx 是什么，可以深度解析下吗？

PTX（Parallel Thread Execution）是由NVIDIA推出的一种GPU程序语言，用于编写在GPU上执行的并行程序。它是一种低级程序语言，类似于汇编语言，但比汇编语言更易于编写和阅读。

img

PTX程序语言的基本单位是线程块（thread block），而线程是在GPU上独立执行的最小单位。线程块内的线程可以通过共享内存进行通信和协作。

PTX提供了一系列基本指令，如加载内存、存储内存、线程同步和分支跳转等。此外，PTX还提供了一些针对GPU架构的优化指令，例如cooperative groups、shared memory和atomic operations等，这些指令能够更好地利用GPU的计算能力。

开发者可以使用NVIDIA提供的CUDA Toolkit进行PTX程序的开发，也可以使用其他第三方工具。在编写PTX程序时，需要考虑GPU的并行架构和线程调度策略，以充分发挥GPU的计算能力。

13. roofline模型有什么用？如何确定最优的BLOCK_SIZE。

Roofline模型是一种用于分析计算平台性能的理论模型，由加州理工大学伯利克实验室提出。该模型基于计算强度（Operational Intensity，OI）和内存带宽两个指标，可以预测不同计算平台在不同计算强度下的理论计算上限。计算强度是指算法的计算量和数据量之比，通常以FLOPS为单位。

在建立Roofline模型时，需要获取计算平台的硬件参数，包括CPU频率、内存带宽、AVX512和FMA等。这些参数可以通过查询CPU手册或使用lscpu等工具获取。

Roofline模型中的理论性能峰值和理论内存带宽可以通过以下公式计算：

理论性能峰值 = 频率 * 512 * AVX数量 * FMA32/64

理论内存带宽 = 内存带宽

计算强度（OI）是指算法的计算量和数据量之比，通常以FLOPS为单位。对于访存密集型算法，可以通过soft prefetch等技术优化；对于计算密集型算法，可以通过SIMD等技术优化。

确定最优的BLOCK_SIZE需要结合具体的硬件环境和程序特点来进行评估。一般来说，较大的BLOCK_SIZE可以提高计算效率，但也会增加内存访问冲突的可能性。因此，在确定最优的BLOCK_SIZE时，需要权衡计算和内存带宽之间的关系，并考虑程序的并行度、数据访问模式等因素。通常可以通过实验和性能分析来寻找最优的BLOCK_SIZE。

14. 稀疏矩阵的存储格式有哪些？稀疏矩阵的应用场景？稀疏矩阵计算与稠密矩阵计算有何不同？

稀疏矩阵的存储格式有以下几种：

  1. COO（Coordinate）格式：将非零元素存储为三元组 (i, j, value)，其中 i 和 j 分别表示元素所在行和列的下标，value 表示元素的值。
  2. CSR（Compressed Sparse Row）格式：将矩阵按行压缩存储为三个数组，分别表示非零元素的值、列下标和每行第一个非零元素在上述两个数组中的位置。
  3. CSC（Compressed Sparse Column）格式：与 CSR 类似，但按列压缩存储。
  4. DIA（Diagonal）格式：将矩阵压缩存储为一个主对角线和若干个副对角线，每个对角线用一个数组存储。

稀疏矩阵计算与稠密矩阵计算有以下不同之处：

  1. 稀疏矩阵计算通常需要使用特定的算法和数据结构，例如稀疏矩阵乘法、最小割算法等。这些算法和数据结构通常需要考虑非零元素的位置和数量，以及存储格式的选择等因素。
  2. 稀疏矩阵计算与稠密矩阵计算在计算复杂度上有所不同。由于稀疏矩阵中大部分元素为零，因此稀疏矩阵乘法等计算的复杂度通常远远低于稠密矩阵乘法。
  3. 稀疏矩阵计算与稠密矩阵计算在硬件实现上也有所不同。例如，在GPU等加速器上，稠密矩阵计算通常可以通过SIMD等技术获得较高的计算效率，而稀疏矩阵计算需要使用特定的算法和数据结构来充分利用硬件资源。

15. 如何计算CPU指令的吞吐量和时延?

在介绍如何计算CPU指令的吞吐量和时延前，我们先来了解下一些基本的概念：

  1. 时钟频率：这是CPU的基本工作频率，以赫兹(Hz)为单位。例如，一个500MHz的CPU，其时钟频率为500,000,000Hz。
  2. 时钟周期：这是CPU完成一个基本操作所需的时间，通常等于时钟频率的倒数。例如，对于500MHz的CPU，其时钟周期为1 / 500,000,000 = 2纳秒。
  3. 机器周期：这是CPU完成一个指令所需的时间，通常包含多个时钟周期。例如，一个指令可能需要3个机器周期才能完成。
  4. 指令吞吐量：这是CPU每秒钟能够执行的指令数，通常以每秒百万条指令(MIPS)为单位。计算公式为：指令吞吐量 = 时钟频率 / 指令周期。
  5. 指令时延：这是CPU执行一个指令所需的时间，通常以时钟周期为单位。计算公式为：指令时延 = 机器周期 * 时钟周期。

例如，假设一个500MHz的CPU，每4个时钟周期组成一个计算机周期，执行一条指令平均需要三个机器周期，那么该处理器的一个机器周期为8纳秒，平均执行速度为41.67 MIPS。

AI 框架知识

1. MLIR有了解过吗？ONNX有了解过吗？

MLIR是一个多级中间表示框架，旨在解决当前深度学习领域中存在的IR（Intermediate Representation，中间表示）碎片化问题。它由LLVM团队开发和维护，强调工具链的可重用性和可扩展性。MLIR试图通过引入统一的表示形式——Dialect（方言），来克服不同IR之间转换难度大、优化Pass难以共享等问题。

MLIR它不仅仅局限于深度学习或某个特定的框架。MLIR旨在提供一个多级抽象的表示，支持从高层次算法到低层硬件指令的各种编程范式和优化技术。MLIR既可以用于深度学习模型的表示和优化，也可以用于传统软件编译领域，甚至是硬件设计的高级表示。

ONNX（Open Neural Network Exchange）是一种开放的IR，主要用于不同深度学习框架之间的模型转换和互操作性。ONNX通过定义一套标准化的算子和模型格式来实现这一点。

无论你使用何种训练框架训练模型（比如TensorFlow/Pytorch/OneFlow/Paddle），在训练完毕后你都可以将这些框架的模型统一转换为ONNX这种统一的格式进行存储。

然而，ONNX主要关注模型的表示和交换，而不涉及模型的优化过程。

2. TVM的整体结构，如何用TVM进行开发？

img

TVM的整体结构可以分为以下几个关键组成部分：

  1. 前端：TVM支持多种深度学习框架的模型作为输入，比如TensorFlow、PyTorch、MXNet、Keras等。通过这些前端接口，TVM可以读取不同框架定义的模型，并将其转换成中间表示（IR）。
  2. 中间表示（IR）：TVM使用两级IR，即Relay和TIR（Tensor IR）。Relay是一种高级IR，用于表示高级神经网络算法；而TIR是一种低级IR，用于表示更接近于硬件的操作和优化。
  3. 自动调度（AutoTVM/AutoScheduler）：为了在特定硬件上获得最佳性能，TVM提供了自动调度工具，如AutoTVM和更现代的AutoScheduler，它们可以自动优化模型的计算图和内核实现。
  4. 运行时：TVM提供了一个轻量级的运行时，支持模型在目标硬件上的部署和执行。这包括对多种设备的支持，如CPU、GPU、FPGA等。
  5. 编译流程：TVM的编译流程包括模型的加载、优化（例如算子融合、内存优化）、自动调度、代码生成等步骤，最终生成可以在目标硬件上运行的机器码。

3. 为什么要进行推理优化？直接用tensorflow或者pytorch的推理接口不行吗？

首先，进行推理优化的原因主要是为了提高模型在实际应用中的性能。

例如，对于电商软件或内容应用实时的个性化推荐来说，它要求能够快速响应，因为推荐的卡顿感将直接影响购物或者内容获取的体验。量化模型毫秒级的交易判断输出能帮助华尔街的交易员们套取巨额利润。

甚至他们愿意牺牲一部分的模型表现来换取更高的推理性能。

其次，随着深度学习的广泛应用和以Transformer大模型为基座的深度算法的普及，tensorflow或者pytorch的推理性能远远跟不上了。

为了提高模型在实际应用中的性能，解决其在推理时面临的高计算复杂度、大内存需求、并行性限制等问题，常用一些办法总结如下：

  1. 模型压缩技术：通过知识蒸馏、权重剪枝、量化等方法，减小模型体积和计算需求。这些技术通过精简模型结构或参数，减少不必要的计算负担，从而加快推理速度，降低存储和运行时内存需求。知识蒸馏特别值得注意，它通过将大模型的"知识"转移到小模型上，既保持了模型性能，又实现了模型大小的显著减少。
  2. 高效变体的开发：针对自注意力机制的高计算复杂度，开发了如Transformers的变体，通过修改自注意力机制来降低计算复杂度和内存需求。例如，使用稀疏注意力模式、局部注意力或低秩近似等技术，有效地减少了处理长序列时的计算负担。
  3. 硬件加速和专用推理引擎：利用TPU、GPU等专门设计的硬件加速器和优化的推理引擎（如TensorRT、ONNX Runtime），针对特定硬件平台的特性进行底层优化。这种优化可以显著提高推理速度，同时降低功耗，特别适合需要实时处理的应用场景。
  4. 动态量化和混合精度推理：通过在模型的不同部分使用不同的数据精度（如FP32、FP16、INT8）来平衡推理速度和模型精度之间的关系。动态量化特别适用于在运行时根据需求调整精度，而混合精度推理则可以在保持模型性能的同时加速模型的推理过程。

4. 模型推理优化的常用方法有哪些？

推理性能优化的常用方法主要包括两大类，一类是模型的压缩技术，比如模型的剪枝、量化、蒸馏等，在较大的预训练模型下被广泛使用；另外是推理加速的技术，在CPU、GPU加速上分别有一些方法。

  1. 网络剪枝

网络剪枝是从大型网络中筛选出不重要的神经元以及权重，将它们从网络中删除，同时尽可能地保留网络的性能。

2. 量化技术

量化技术是把高精度表示的网络权重和激活值，用低精度来近似表示，实现网络的轻量化。优势如下:

网络存储 ：每个层权重量化后，32位的比特就可以压缩到8比特，就是浮点型到整形的量化，整个模型占的空间就会变小；

激活值 : 通过使用较少位的数值表示，在处理同样数据时需要读/写的内容就更短，内存带宽的压力就变得更小；

计算时间 ：单位时间内处理定点运算指令就会比浮点运算的指令多。

3. 模型蒸馏

蒸馏是一种模型压缩常见方法，将复杂、学习能力强的网络学到的特征，表示“知识”蒸馏出来，传递给参数量小、学习能力弱的网络。

img

知识蒸馏需要两种类型的网络：Teacher Model和Student Model。前者参数量大、结构复杂，后者参数量较小、结构相对简单。二者可以是不同的网络结构，但是采用相似的网络结构，蒸馏效果会更好。

4. Caching

CPU频率远快于主存访问速度，在处理器时钟周期内，CPU常常需要等待主存，浪费计算资源。为了缓解CPU和内存之间速度的不匹配问题，增加了CPU cache 来解决。

cache 利用局部性原理来提高缓存命中率:

A. 时间局部性：如果某个数据被访问，那么在不久的将来它很可能被再次访问；

B. 空间局部性：如果某个数据被访问，那么与它相邻的数据很快也可能被访问。

5. 多算子融合

img

算子融合是GPU上一个很重要的推理加速的一个优化手段，尤其是针对NLP这样的大模型，会带来比较显著的效果的提升。对于GPU异构编程，每一次op操作都会有一个内核的调用和多次的显存的读取；对于小op来说启动GPU kernel的时间会大于GPU计算时间，显存的读取开销也很大；op数目太多的话,效率会变低；所以将算子合并，可以有效地提高计算的性能。

6. 计算图优化

每个计算图中都包含许多计算节，图优化的目标很简单，就是简化计算图中计算节点的计算量。常用的方式分为以下几种：

  1. 减少节点的数量
  2. 用高效替换低效的节点
  3. 用高效子图替换低效子图
  4. 用并行化分支代替单分支

5. TensorRT如何进行自定义算子开发？

首先，明确为什么需要自定义算子，了解TensorRT中自定义算子的基本概念，包括插件（Plugin）和插件工厂（Plugin Factory）。

其次，开始实现自定义算子，主要分为三步，定义插件类，实现算子逻辑和注册插件，下面分别展开说说。

  1. 定义插件类
  • 定义插件类：继承nvinfer1::IPluginV2接口，实现其虚函数，包括算子计算（enqueue）、输出维度（getOutputDimensions）、数据类型和格式配置（configureWithFormat）、初始化与清理（initialize 和 terminate）、序列化与反序列化（serialize 和 deserialize）等。
  1. 实现算子逻辑
  • 核心实现：在enqueue函数中实现算子的具体计算逻辑，可能需要使用CUDA等技术在GPU上进行并行计算以提高效率。
  1. 注册插件
  • 注册插件：通过实现nvinfer1::IPluginCreator接口并注册插件，使得TensorRT能够识别和使用自定义算子。

再次，是构建使用自定义算子代码

  • 创建插件实例：在网络构建过程中，使用插件工厂或直接调用插件类的构造函数创建自定义算子实例，并将其嵌入到模型中。
  • 序列化和反序列化：确保自定义算子能够被正确地序列化和反序列化，以便模型的保存和加载。

最后，测试和验证

  • 验证正确性和性能：通过与预期结果或其他框架的对比，验证自定义算子的实现是否正确、性能是否达标，并确保其在不同条件下的鲁棒性。

img

6. TensorRT对模型实现了哪些推理优化？

  1. 常量折叠（Constant Folding）
  • 解释：常量折叠是一种编译时优化技术，它预计算图中那些在推理前就能确定结果的表达式。这意味着网络中的任何常量操作，如常量之间的算术运算，都会在模型编译期间被提前计算并简化，从而减少运行时的计算负担。
  • 好处：通过减少不必要的运行时计算，可以显著提高模型的推理速度。
  1. 算子融合（Layer/Operator Fusion）
  • 解释：算子融合是将多个操作合并为一个复合操作的过程，这可以减少内存访问次数并降低推理延迟。例如，卷积、批量归一化（Batch Normalization）、激活函数等可以融合成一个单一的高效操作。
  • 好处：减少了内存访问和计算步骤，提高了数据吞吐率和运算效率。
  1. 量化（Quantization）
  • 解释：量化是指将模型中的权重和激活从浮点数转换为低精度的表示形式，如从32位的浮点数（FP32）转换为8位整数（INT8）。TensorRT提供了量化校准工具，以最小化量化带来的精度损失。
  • 好处：量化可以显著减少模型的大小和推理时间，同时也减少了内存带宽的需求，使得模型更适合在资源受限的设备上运行。
  1. 层自动调整（Layer Auto-Tuning）
  • 解释：TensorRT会针对特定的硬件平台自动选择最优的算法来执行各种操作。这是通过在不同的算法实现之间运行基准测试来完成的，以确保选择的实现能够提供最佳性能。
  • 好处：确保模型在特定硬件上达到最优的运行效率。
  1. 多流执行（Multi-Stream Execution）
  • 解释：TensorRT支持利用GPU的并行处理能力，通过多流执行来同时处理多个推理请求，从而提高吞吐量。
  • 好处：在处理大量并发请求时，能够有效提升GPU利用率和总体吞吐量。

7. 算子融合为什么能加速推理，优化了哪一部分？TensorRT用到了哪些算子融合？算子融合在推理框架中是如何实现的？

首先，算子融合通过将多个操作合并为一个单一复合操作来减少模型中的层次。这种优化减少了内存访问次数和数据传输量，因为它降低了中间结果的读写需求。同时，减少了CPU和GPU之间的同步点，提高了计算效率。

其次，TensorRT利用算子融合技术，例如将卷积、批量归一化（Batch Normalization）、激活函数（如ReLU）融合为一个单一操作。这种融合不仅减少了计算步骤，还减少了中间数据的存储需求。

img

在推理框架中，算子融合通常在图优化阶段进行。框架会分析计算图，识别可以融合的操作序列。然后，通过生成一个包含所有融合操作逻辑的新内核来实现融合。这个过程可能涉及到自动生成代码或者使用预先定义的高效内核。

然而，实现算子融合需要考虑操作之间的依赖关系和数据流动，确保融合后的操作不会引入错误。此外，还需要平衡融合的程度和实现的复杂性，以达到最佳性能。

8. 模型量化的加速原理，模型量化带来的精度损失如何解决？

首先，模型量化是指将模型中的权重和激活函数的数据类型从浮点数（如FP32）转换为低精度的格式（如INT8或FP16）。

这种转换减少了模型的内存占用，降低了计算所需的内存带宽，并且在一些硬件上可以利用专门的低精度计算单元，从而加速模型的推理过程。

量化可以显著减少模型大小，提高数据加载和处理速度。例如，在支持INT8指令集的硬件上，使用INT8量化的模型相比于FP32模型，理论上可以提高4倍的内存带宽效率和计算速度。

量化带来的精度损失及解决方案

精度损失原因：量化过程中，由于将浮点数映射到有限的整数范围，会引入量化误差，导致模型的精度下降。

解决策略：

  1. 量化校准：通过在量化前后对模型进行校准，选择最优的量化参数（如量化比例和零点），以最小化量化误差。常用的校准方法包括最小最大值校准、百分位校准等。
  2. 量化感知训练：在模型训练过程中模拟量化的效果，让模型“适应”量化带来的误差。这种方法可以在训练阶段就考虑到量化误差，进而学习到更鲁棒的权重。

模型量化是一种有效的推理加速技术，虽然可能会带来一定的精度损失，但通过量化校准和量化感知训练等策略，可以显著减轻甚至克服这一问题。

9. ONNX Runtime支持在多种硬件上进行推理，说明具体的实现机制。

首先，ONNX Runtime是一个用于优化和运行机器学习模型的性能引擎。它支持使用ONNX格式的模型，这是一个开放格式，用于表示机器学习模型，使得不同的AI框架训练的模型能够在不同的平台和设备上运行。

其次，多硬件支持的实现机制实现原理是，ONNX Runtime通过提供一系列的“执行提供程序”（Execution Providers，EPs）来支持不同的硬件。每个执行提供程序都是为特定的硬件或计算库优化的后端，它定义了如何在该硬件上执行ONNX模型中的操作。它会根据可用的执行提供程序和硬件资源自动选择最适合当前环境的执行路径。用户也可以手动指定使用哪个执行提供程序，以便更精细地控制模型的运行方式。

img

例如，NVIDIA GPU的系统上，ONNX Runtime可以利用CUDA执行提供程序来加速模型的推理。此外，对于需要进一步优化的场景，可以使用TensorRT执行提供程序，它利用NVIDIA TensorRT进行图优化和内核融合，以实现更高效的推理。

总之，ONNX Runtime通过灵活的执行提供程序机制，有效地支持了多种硬件平台，这种跨平台的能力大大降低了模型部署的复杂性，并为开发者提供了更多的灵活性和选择。

10. 总结一下TensorRT，ONNX Runtime等推理框架的组成架构，如果我们公司自己要为硬件开发一套推理框架，应该重点关注哪些部分？

从TensorRT，ONNX Runtime推理框架来看，推理框架的核心组成部分应该分为以下几个部分

  1. 模型解析器：负责将机器学习模型（如ONNX格式）转换成框架能够理解和执行的内部格式。
  2. 图优化器：通过算子融合、常数折叠等技术优化计算图，减少不必要的计算，提高执行效率。
  3. 执行引擎：负责根据优化后的计算图，在特定硬件上执行计算任务。
  4. 硬件抽象层（对应于TensorRT、ONNX Runtime中的执行提供者/后端）：为不同硬件提供定制化支持，确保模型能够在多种平台上高效运行。

在为硬件开发一套推理框架时，重点要关注两个方面，兼容和利用硬件特点优化性能和图优化技术。

推理框架需要根据目标硬件的特点进行专门的优化，通过利用硬件的并行计算能力、特殊的指令集、高效的内存访问，显著提高模型的推理速度，降低能耗，突出差异化和优势。

通过算子融合、优化数据传输路径等，减少不必要的计算和内存访问，提升模型执行的效率。这对于加速模型推理、减少资源消耗至关重要。

11. 各种推理框架都有何优劣势？它们的性能怎么样？

img

  1. TensorRT
  • 优势：专为NVIDIA GPU设计，提供高度优化的推理性能，特别适用于高吞吐量的服务器端和边缘设备场景。支持精确的算子融合、量化、动态张量等高级优化。
  • 劣势：主要限制在NVIDIA GPU上，不适用于其他类型的硬件。相对其他框架，学习曲线可能更陡峭。
  • 性能：在NVIDIA GPU上，TensorRT通常能提供最佳的推理速度和效率。
  1. ONNX Runtime
  • 优势：支持多种硬件平台，包括CPU、GPU和FPGA等。与ONNX模型格式紧密集成，方便从不同的训练框架迁移模型。微软背书，社区活跃。
  • 劣势：虽然支持多种硬件，但在特定硬件上的优化可能不如专门的推理引擎深入。
  • 性能：提供良好的跨平台性能，但在特定硬件上可能不是最优。
  1. TensorFlow Lite
  • 优势：专为移动和嵌入式设备优化，支持模型量化和轻量级操作，减少模型大小和提升推理速度。广泛应用于Android和iOS设备。
  • 劣势：相比于服务器端的框架，可能在功能和性能上有所限制。
  • 性能：在移动设备上提供优化的推理性能，特别是在支持NEON指令集和GPU加速的设备上。
  1. Core ML
  • 优势：苹果官方支持，为iOS、macOS、watchOS和tvOS设备上的机器学习应用优化。可以直接利用苹果设备的硬件加速能力。
  • 劣势：限于苹果生态系统，不适用于非苹果平台。
  • 性能：在苹果设备上提供高效的推理性能，特别是通过使用Metal进行GPU加速时。

性能总结

推理框架的性能受到多种因素影响，包括模型的复杂度、硬件的计算能力、框架的优化程度等。一般而言，专门为特定硬件优化的推理框架（如TensorRT、Core ML）能够提供最佳性能。然而，跨平台框架（如ONNX Runtime、TensorFlow Lite）提供了更广泛的适用性和灵活性，允许模型在多种设备上运行，虽然可能牺牲一定的性能。

12. 分布式训练中有哪些并行模式？每种模式需要做什么，有什么优缺点？

分布式训练主要分为两种数据并行和模型并行。

数据并行适用于大规模训练数据集和相对较小的模型，能够提高训练速度。通常需要将训练数据集划分为多个子集，每个设备或节点负责处理一个子集，并在每个子集上独立训练模型。然后，通过梯度聚合和同步来更新模型参数。

它的缺点是需要大量的通信开销，因为设备之间需要传输梯度信息。同时，需要额外的内存来存储模型副本和梯度。

模型并行适用于大型模型或需要更高计算需求的任务，允许训练更大规模的模型。它指的是将模型分解为多个部分，在不同的设备或节点上并行处理。每个设备只负责处理模型的一部分，并与其他设备交换中间结果。

缺点是需要更复杂的编程和通信模式，以确保各个设备之间的协同工作。可能存在设备之间的通信瓶颈。

不过在实际应用中，一般是可以将数据并行和模型并行结合使用，以充分利用多个设备和节点的计算能力。

13. 分布式训练中我们重点需要处理的问题有哪些？

  1. 通信开销和延迟：

在分布式环境下，非常需要关注设备之间的通信开销和传输延迟。为了减少通信开销，可以采用压缩、稀疏化和量化等技术。同时，可以利用高速网络和专用硬件来减少传输延迟。

例如，可以使用混合精度训练（Mixed Precision Training）将参数和梯度从单精度浮点数压缩为半精度浮点数，从而减少传输数据的大小。在多机之间进行数据传输时，可以采用RDMA（Remote Direct Memory Access）网络或InfiniBand网络等减少传输延迟。

2. 容错性和可伸缩性

在分布式环境下，尤其是多机训练中，必须考虑的就是训练中的设备故障或网络异常。例如断点训练，冗余参数服务器等。

14. MPI可以应用于AI框架的哪些方面？

  1. 数据并行性：
  • AI模型的训练通常需要处理大规模的数据集，而数据并行性可以将数据划分为多个部分，并将其分发到不同的计算节点上并行处理。MPI提供了消息传递的能力，可以在不同计算节点之间传输数据，实现数据的并行处理。

2. 参数同步：

    • 在分布式训练中，不同计算节点上的模型参数需要进行同步，以保持一致性和提高训练效果。MPI可以通过消息传递的方式，在计算节点之间进行参数的同步和更新，确保所有节点上的模型参数保持一致。

3. 集群扩展：

    • MPI可以应用于AI框架中的集群扩展，使得框架可以在大规模分布式环境下运行。通过MPI，可以实现节点之间的通信和数据传输，以及任务的分发和调度，从而实现在集群中进行规模化的训练和推理。

4. 容错性：

    • AI框架中的分布式训练通常需要考虑设备故障和恢复的情况。MPI提供了容错功能，可以在计算节点故障后重新分配任务和数据，并进行恢复，以保证训练过程的连续性和可靠性。

可以参考Horovod在TensorFlow、PyTorch深度学习训练中的使用。

15. 反向传播的原理，具体实现的源码有了解过吗？

深度神经网络的权重是如何精确调整的？ 它们就是通过 反向传播 进行调整的。 如果没有反向传播，深度神经网络将无法执行识别图像和解释自然语言等任务。

反向传播的目标 是为了减少损失 或误差，通过调整网络的权重来实现的，使假设更像输入特征之间的真实关系。

前向传播是神经网络中的常规训练过程，通过对输入进行加权求和和激活函数运算得到输出，然后将该输出作为下一层的输入，一直传递到神经网络的末端。

反向传播过程接受模型训练过程的最终决策，然后确定这些决策中的错误。 通过对比网络的输出/决策和网络的预期/期望输出来计算误差。

一旦计算出网络决策中的错误，该信息就会通过网络反向传播，并且网络参数会随之改变。 用于更新网络权重的方法基于微积分，具体来说，它基于链式法则。

当神经元提供输出值时，会使用传递函数计算输出值的斜率，从而产生派生输出。 进行反向传播时，特定神经元的误差根据以下公式计算 公式 :

误差 = (预期输出 – 实际输出) * 神经元输出值的斜率

对输出层的神经元进行操作时，将类别值作为期望值。 计算出误差后，该误差将用作隐藏层中神经元的输入，这意味着该隐藏层的误差是输出层中找到的神经元的加权误差。 误差计算沿着权重网络向后传播。

计算出网络的误差后，必须更新网络中的权重。 如前所述，计算误差涉及确定输出值的斜率。 计算出斜率后，将进行一个称为 梯度下降 可用于调整网络中的权重。

梯度是一个斜率，其角度/陡度可以测量。 斜率是通过在“运行”上绘制“y”或“上升”来计算的。 在神经网络和错误率的情况下，“y”是计算的误差，而“x”是网络的参数。 网络的参数与计算的误差值有关，并且随着网络权重的调整，误差会增加或减少。

“梯度下降”是更新权重以降低错误率的过程。 反向传播用于预测神经网络参数与错误率之间的关系，从而建立梯度下降网络。 使用梯度下降训练网络涉及通过前向传播计算权重、反向传播误差，然后更新网络的权重。

16. TensorFlow和Pytorch都用过吗？它们设计思路有何不同？有何优劣？如何添加自定义算子？

TensorFlow像是搞数据的人研发的，Pytorch才像是搞算法人研发的。

Tensorflow 和 pytorch 相比，前者占据先机，后者则势头完全盖过前者。

TensorFlow：由Google开发，设计时考虑了分布式计算、大规模数据处理和生产环境的需求。TensorFlow使用静态计算图，需要先定义后运行，适合于大型模型和复杂的神经网络。

PyTorch：由Facebook的人工智能 研究团队开发，设计理念侧重于灵活性和直观性。PyTorch采用动态计算图，允许即时修改和执行，非常适合于快速原型开发和研究工作。

TensorFlow：虽然提供了强大的功能，但其静态图机制使得调试和理解模型相对复杂。TensorFlow 2.0及后续版本引入了Eager Execution，改善了易用性。但API得兼容性也是一个大问题。

PyTorch：动态图的特性使得PyTorch在模型构建和调试方面更加直观和用户友好。PyTorch的API设计更接近Python原生，易于理解和使用。

TensorFlow：在大规模数据集和复杂模型的训练上展现出较强的性能，尤其是在GPU加速和TPU支持方面。

PyTorch：虽然在早期版本中性能略逊于TensorFlow，但近年来通过优化和社区努力，其性能已大幅提升，尤其在某些特定任务上表现出色。

TensorFlow：由于其稳定性和规模化部署能力，非常适合用于商业产品和大型企业项目。TensorFlow Serving、TensorFlow Lite和TFX等工具支持了从研发到生产的全流程。

PyTorch：以其灵活性和友好的API，更受研究人员和数据科学家的喜爱，特别是在进行快速实验和研究原型开发时。不过随着大语言模型的爆火，很多大语言模型的商用都是基于PyTorch。

img

img

在TensorFlow和PyTorch中，可以通过以下步骤来添加自定义算子：

TensorFlow:

首先，是实现底层算子，可以使用C++或者CUDA，来写出你想要实现的功能，也就是自定义算子的代码。把这段代码编译成一个库文件，这样我们的TensorFlow才能认识它。

其次，需要在TensorFlow里面创建一个新的操作符（Op），并通过注册机制告诉TensorFlow，这时需要提供这个算子的名字，输入输出是什么类型的，形状长啥样等等。

最后，在Python环境下，你可以用TensorFlow提供的API来把这个自定义算子加入到你的计算图中。通过tf.load_op_library()函数加载你的库文件，然后用tf.custom_op()函数把它当作一个操作来使用。

PyTorch:

对于PyTorch，其实步骤差不多，但是在PyTorch中你可以直接用Python或者C++来写你的自定义算子代码。 然后，在Python环境下，你利用torch.utils.cpp_extension来编译这段代码，生成一个扩展模块。这里同样需要告诉它算子的名字，数据类型和形状等信息。 编译好之后，就可以像导入普通模块一样，把这个扩展模块导入到PyTorch中了。

最后，使用这个自定义算子就像使用PyTorch中其他的操作一样简单。

无论是在TensorFlow还是PyTorch中添加自定义算子，过程都是：“写算子，编译，注册，使用”。

  • Title: CUDA-interview-note
  • Author: Charles
  • Created at : 2024-04-15 18:10:17
  • Updated at : 2024-04-15 18:28:09
  • Link: https://charles2530.github.io/2024/04/15/cuda-interview-note/
  • License: This work is licensed under CC BY-NC-SA 4.0.
recommend_articles
computer-network-note-4 computer-network-note-4 OO_multithreading_tips OO_multithreading_tips OO_pre_7 OO_pre_7
recommend_articles
computer-network-note-4 computer-network-note-4 OO_multithreading_tips OO_multithreading_tips
Comments
On this page
CUDA-interview-note
  1. 整体情况简介
  2. 高性能计算基础
  3. AI 框架知识
  4. 算法题
  5. 一些比较零碎的问题
  6. 推荐参考资料
  7. 高性能计算基础
    1. 1. 你可以简单说下CUDA的线程组织结构吗？
    2. 2. 可以简单谈谈CUDA的存储体系结构，每一种存储的优缺点，该如何合理使用。
    3. 3. 你了解CUDA stream吗？为什么要使用多个stream？
    4. 4. GPU和CPU分别适合执行哪些程序？结合它们的硬件架构解释一下为什么它们有各自的优势。
    5. 5. 说明一下神经网络加速器与CPU、GPU的区别，他们各自有何优势？
    6. 6. 半精度浮点数FP16各个部分的具体位数，为什么要有半精度浮点数？
    7. 7. 可以谈下TensorCore的加速原理吗？
    8. 8. MPI，OpenMP以及CUDA各自适用的加速场景。
    9. 9. 可以说下DMA和RDMA是什么吗？以及有哪些硬件上的实现。
    10. 10. 平时如何进行kernel的优化，会用到哪些工具？
    11. 11. CPU上哪些并行优化方法？
    12. 12. ptx 是什么，可以深度解析下吗？
    13. 13. roofline模型有什么用？如何确定最优的BLOCK_SIZE。
    14. 14. 稀疏矩阵的存储格式有哪些？稀疏矩阵的应用场景？稀疏矩阵计算与稠密矩阵计算有何不同？
    15. 15. 如何计算CPU指令的吞吐量和时延?
  8. AI 框架知识
    1. 1. MLIR有了解过吗？ONNX有了解过吗？
    2. 2. TVM的整体结构，如何用TVM进行开发？
    3. 3. 为什么要进行推理优化？直接用tensorflow或者pytorch的推理接口不行吗？
    4. 4. 模型推理优化的常用方法有哪些？
    5. 5. TensorRT如何进行自定义算子开发？
    6. 6. TensorRT对模型实现了哪些推理优化？
    7. 7. 算子融合为什么能加速推理，优化了哪一部分？TensorRT用到了哪些算子融合？算子融合在推理框架中是如何实现的？
    8. 8. 模型量化的加速原理，模型量化带来的精度损失如何解决？
    9. 9. ONNX Runtime支持在多种硬件上进行推理，说明具体的实现机制。
    10. 10. 总结一下TensorRT，ONNX Runtime等推理框架的组成架构，如果我们公司自己要为硬件开发一套推理框架，应该重点关注哪些部分？
    11. 11. 各种推理框架都有何优劣势？它们的性能怎么样？
    12. 12. 分布式训练中有哪些并行模式？每种模式需要做什么，有什么优缺点？
    13. 13. 分布式训练中我们重点需要处理的问题有哪些？
    14. 14. MPI可以应用于AI框架的哪些方面？
    15. 15. 反向传播的原理，具体实现的源码有了解过吗？
    16. 16. TensorFlow和Pytorch都用过吗？它们设计思路有何不同？有何优劣？如何添加自定义算子？