·设为首页收藏本站📧邮箱修改🎁免费下载专区💎积分✅卡密📒收藏夹👽聊天室
DZ插件网 门户 站长资讯 查看内容

AI深度学习编译器工程师需要哪些技术栈?

2023-8-5 00:16| 发布者: IT618发布| 查看: 22353| 评论: 0

摘要: 了解AI深度学习编译器工程师的技术栈,包括深度学习框架、编译原理、算子编写、硬件工程师和高性能工程师等知识。通过抓住重点、构建框架和灵活深入的学习方法,积累解决问题的能力。同时,掌握入门重点、框架和快速深入的方法,以提高针对具体工作需求的知识和能力。
AI深度学习编译器工程师需要哪些技术栈?7794 作者: 来源: 发布时间:2023-8-5 00:16

AI深度学习编译器工程师需要哪些技术栈?5594 作者: 来源: 发布时间:2023-8-5 00:16



AI深度学习编译器工程师需要哪些技术栈?6726 作者: 来源: 发布时间:2023-8-5 00:16
https://www.zhihu.com/question/532768471/answer/2692111925

大概两年半前入了这个坑,就一直在思考这个问题。

深度学习编译器需要了解的知识面大的恐怖,理想的情况是:
  • 像算法工程师一样了解各个模型的结构
  • 像深度学习框架工程师一样了解各个功能模块的实现
  • 像编译器工程师一样了解编译原理
  • 像高性能工程师一样了解算子编写及优化方法
  • 像硬件工程师一样了解体系结构
  • 还有基础的编程语言,数据结构算法,操作系统原理....

这显然不现实,你不可能在每个子领域都和专注这个子领域的人懂得一样多,这也是困扰我挺长一段时间的问题,学的再多,似乎也只是沧海一粟,这个清单要仔细地列出来,远比楼上的高赞回答长得多。

于是我只能退一步想,技术栈的目的是为了解决问题,那么问题就可以变成,如何分配有限的时间,最大化积累的解决问题能力。找到最有学习价值的内容,也许比学习本身更重要。知道这么是暂时不需要了解的,比知道什么是需要了解的更宝贵。

具体来说,我把我的方法大致总结为“抓住重点,构建框架,灵活深入”。抓住重点,指的是每个领域都有一些入门的,必须掌握的,不懂就完全看不了其他的知识。构建框架,指的是在了解了基础之上,对知识有框架性的认识,知道这个领域要解决哪些问题,哪些问题大致有哪些方法,哪些方法可以去哪里看到。灵活深入,以前两点作为基础,也是最后要达到的效果,既然我们难以做到直接充分掌握各个领域的技术细节,那我们就应该追求,当工作中分配到某个任务的时候,基于任务的需求,快速地获得这项任务需要的知识。

先挖个坑,回头再写各个领域具体哪些是入门重点,框架如何,快速深入的方法是什么。

这里插播一个声明,我说的这些主要是针对真正的具体工作需要积累哪些知识和能力。至于面试会问啥,还是多搜搜面经吧,毕竟国内技术面试脱离实践这个问题也不是一两天了。

先填一下算法层的坑吧。算法工程师要解决的问题是,端到端地解决具体的问题,具体来说就是从用户给出的原始数据到用户需要拿到的最终数据,比如如果是一个对图像的分割任务,那就是从原始的图片数据集到最后标号分割框的输出图片。而对于深度学习编译器来说,要解决的问题是对于一个已知结构的模型,找到最优的计算实现,这个链路,是比算法工程师要短很多的。更进一步来说,深度学习编译器只需要理解一个模型长什么样。至于模型有什么用,模型为什么长这样,怎么训练出这个模型,怎么获得这个模型想要的输入,并不重要。

举例来说,算法工程师要知道怎么处理训练数据不平衡的问题,这跟深度学习编译器就没什么关系。又比如NLP里面,输入数据要怎么做embedding,也不重要,又比如,bert在最后一层加上不同的尾处理,就可以拿来解决不同的问题,这也不重要。

那理解一个模型需要长什么样,具体需要懂啥呢?最基本的当然是各个算子的语义,这个没什么好说的,应该都能掌握,对自己要求高一点的话,可以对着ONNX文档一个个看过去。至于模型层面,模型的数量是浩如烟海的,但是常见的模块来来回回就是那些,深度学习编译器的很多优化都是基于这些模块来的。所以我会建议大家把重点放在积累理解常见的模型模块,比如conv-bias-batchNorm,比如transformer的encoder/decoder blocker。对于整个模型来说,不需要看太多了,看一下最经典的几个网络,也就是MLPerf跑的那几个就好了,resnet,yolov3,bert。

在工作的时候,你完全可能会遇到一个相对陌生的模型,所以你需要具备快速把握一个模型的能力,这里推荐大家积累的具体能力有两个。一个是通过netron阅读模型的onnx能力。在你掌握上一段说的理论内容之后,你可以下载一个模型的onnx文件试着打开看看,刚开始也许还是蛮痛苦的,看过的应该都懂,但你应该积累出的能力是,打开一个陌生模型的onnx文件,很快能够看出这个模型大致包括哪些基本结构模块,进而理解优化要点有哪些。第二个能力是阅读模型代码能力,比如NLP方面的HuggingFace, CV方面的OpenMMlab,有些时候你需要了解一些模型细节,直接进到代码里面看是更方便的。

今天来填填体系结构的坑,目前还不熟悉NV生态之外的硬件,就以NV作为代表来说了,以及,我会把对CUDA的掌握放在这个部分。这个技术点的重要性取决于你是否希望走性能优化这条路,要理解GPU的性能,就必须深刻地理解体系结构。计算机里面讲的体系结构,我理解就是,对于软件设计有帮助的,不涉及具体电路设计的所有硬件细节。在NV的生态里,CUDA其实就是GPU体系结构的软件抽象。所以,了解CUDA是了解GPU的窗口,而进一步了解GPU的体系结构,其实就是在更深刻地理解CUDA。进一步说,认知建立的步骤应该是,先从功能性的角度过一遍CUDA,有一个整体性的,框架性的认知;然后学习GPU的体系结构,能否充分地发挥硬件性能,本质上在于能否充分利用硬件的并发性和局部性,所以对GPU体系结构的掌握,需要知道每一个可以提高并发性和局部性的硬件特性;最后一个步骤是带着对硬件的理解,再回到CUDA,理解如何通过CUDA来利用各种硬件特性,或者说规避各种性能陷阱。这里推荐三个学习资源,分别对应上述说的三个步骤,第一个是NV官方的CUDA Sample, 第二个是景乃峰,柯晶和梁晓峣的《通用图形处理器设计》,第三个是官方的CUDA Programming Guide。如果还想继续深入的话,可以继续学习PTX,如果还想学习指令集(SASS),也许可以看一下开源的逆向工程做的汇编器,比如商汤最近开源的CUAssmbler。

今天来填一下高性能计算库的坑,所谓高性能计算库,指的是面向一类计算问题的kernel库,比如cublas, cutlass, cutensor, cub(这里我没有提cudnn,cudnn最近几个版本的演化已经有点从算子库向一个处理局部图的框架走的感觉了)。开发高性能算子库大概需要三层能力。

第一层是如何针对特定的问题写出高效的kernel,以GEMM为例,不同的problem size(M,N,K),数据类型,layout在不同的卡上遇到的瓶颈,和对应的最优实现方式都是不一样的。所以你首先得掌握使用各种工具(比如nsight compute)进行性能分析,找到性能瓶颈的方法。然后你得知道算子实现的各种“玩法”,比如各种流水编排,各种优化访存pattern的手段。最后整个流程积累下来,你是应该对这个算子有性能建模的能力的。

第二层是让算子库变得模块化和工程化,一个算子库是面向一类计算问题的软件,不是简单的kernel的集合,作为一个成熟的软件工程,高性能算子库需要尽可能将各种“玩法”变成像积木一样可以自由拼接的模块,最大程度地实现代码复用,同时也方便对整个代码库的管理。除此之外,各种activation操作,add bias,sclase(可以统称epilogue)也需要模块化,并允许被fused到gemm/conv的kernel里面,这对上层的图优化是一个非常重要的支撑。再进一步,代码是死的,团队是活的,从算子库团队而言,应该是一个高性能算子的高效的工厂,新的需求转化成新的功能,或者新的kernel,新的kernel投入测试,解决回归测试中发现的性能问题,每一环的效率都有很大的讲究。

第三层是帮助用户解决算子选择的问题,这个问题有点像第二层做完之后又回到第一层的问题,就是面对一个具体的用户case,怎么找到那个最优的实现。这一层其实已经算是高性能算子库和框架/DL编译器的交界问题了,有些算子库,比如cublas,会有heuristic来一定程度上解决这个问题,而cutlass相比之下就是一个更干净的模板库,并不打算解决这个问题。

让我们回到深度学习编译器研发工程师的视角。深度学习编译器和高性能算子库有两种关系:

第一种关系是深度学习编译器本身就是高性能更算子库,深度学习编译器往往会有自动生成高性能算子的模块(codegen),或者有一些自己写的高性能kernel,如果你是负责这部分的工程师,那你应该就得完全具备高性能算子库工程师的能力了,不过目前深度学习编译器的codegen做得比较好的基本上是pointwise这类实现起来相对简单的算子。

第二种关系是深度学习编译器会调用外部的高性能算子库,这也主要是面对gemm/conv这些计算密集型,玩法比较多,难度比较大的算子。负责管理对外部库调用的工程师就可以更有针对性地积累第二层和第三层的技术了,对于第一层来说,可以暂时只需要了解有哪些“玩法”,每种“玩法”是为了解决什么问题,对于具体的实现方法,优先级可以稍稍往后放。


AI深度学习编译器工程师需要哪些技术栈?7785 作者: 来源: 发布时间:2023-8-5 00:16
https://www.zhihu.com/question/532768471/answer/2684278494

勉强把自己定义为相关人员,以个人工作情况来看,从上到下可能需要以下知识点:
  1. 常见的深度学习算法,检测,识别,分割等等的基本原理;
  2. 常见的模型结构比如卷积,池化,激活函数。典型模型比如mobile net,resnet;
  3. 基本的模型训练流程,部分模型训练技巧,模型结构调整,模型的参数固化和导出;
  4. 在原理层面或者算法层面的模型量化工作;
  5. 常见模型存储格式的解析,例如onnx的模型结构,参数信息解析;
  6. 能自定义简单的模型ir,保证可以比较正确的存储和传递模型信息;
  7. 能理解硬件的设计,知道设计原理更佳;
  8. 可以根据硬件限制做模型结构的融合,拆分,转换;
  9. 能手写常见算子的定点计算(功能层面);
  10. 能理解硬件定点计算方法;
  11. 知道基本的优化器方案,可以比较宏观的理解模型计算流水过程;
  12. 知道dfs,动态规划等基础算法;
  13. 知道二维层面的内存分配,计算时序原理;
  14. 知道cost model概念并可以对不同算子实现不同情况的cost model计算;
  15. 能理解硬件指令,寄存器配置原则,调度原则;
  16. 可以理解基本的运行时工作流程;
  17. 能及时友好沟通简单总结了一部分知识点,如需面试可以按此准备,至少可以进二面。


AI深度学习编译器工程师需要哪些技术栈?6455 作者: 来源: 发布时间:2023-8-5 00:16
https://www.zhihu.com/question/532768471/answer/2778924148

纸上得来终觉浅,绝知此事要躬行。参考TVM文档向Relay添加新算子( TVM对应文档:https://tvm.apache.org/docs/dev/how_to/relay_add_op.html),文档中省去了许多细节,在本文中将实践向Relay添加新算子,并将细节一一列出,给出涉及的文件地址、新增代码、测试脚本。

算子定义


LayerNorm由于其推理时在线计算均值和方差的特性,使得其运行时开销较大LayerNorm计算特性和部署优化。为了减小开销,其中一种方法是采用新的归一化方法替代LayerNorm。RMSNorm就是一个可行的研究工作。

RMSNorm论文:Zhang B, Sennrich R. Root mean square layer normalization[J]. Advances in Neural Information Processing Systems, 2019, 32

对LayerNorm成功的一个著名的解释是它的重新定心和重新缩放的不变性。前者使模型对输入和权值上的移位噪声不敏感,当输入和权值都被随机缩放时,后者保持输出表示的完整。RMSNorm论文假设重新缩放不变性是LayerNorm成功的原因,而不是重新定心不变性。

RMSNorm只关注重新缩放不变性,并根据均方根(RMS)统计量来归一化。相对于LayerNorm,删除了关于均值的统计。

RMSNorm计算公式如下:

AI深度学习编译器工程师需要哪些技术栈?1808 作者: 来源: 发布时间:2023-8-5 00:16

定义属性节点


属性是编译时已知的固定参数。定义一个属性结构体来描述算子的属性。例如Conv2d算子,stride、padding、dilation、kernel_size等为其属性

在tvm/include/tvm/relay/attrs/nn.h添加以下代码:
/*! \brief Attributes used in RMSNorm operator */
struct RMSNormAttrs : public tvm::AttrsNode<RMSNormAttrs> {
  int axis;
  double epsilon;
  bool scale;

  TVM_DECLARE_ATTRS(RMSNormAttrs, "relay.attrs.RMSNormAttrs") {
    TVM_ATTR_FIELD(axis).set_default(-1).describe("Specify which shape axis denotes the channel.");
    TVM_ATTR_FIELD(epsilon).set_default(1e-5).describe(
        "Small float added to variance to avoid dividing by zero");
    TVM_ATTR_FIELD(scale).set_default(true).describe(
        "If true, multiply by gamma; otherwise, gamma is ignored.");
  }
};  // struct RMSNormAttrs

编写类型关系


在编译时需要对算子的输入、输出的类型进行检查,并对算子的输入、输出类型之间的关系进行类型化。这些关系被表示为函数,它接收一个输入类型和输出类型的列表(这些类型中的任何一个都可能是不完整的),并返回一个满足关系的输入和输出类型的列表。这包括形状信息,可以在编译时静态地确定。

在tvm/src/relay/op/nn/nn.cc添加以下代码:
// rms_norm  注册属性节点
TVM_REGISTER_NODE_TYPE(RMSNormAttrs);

//类型检查 形状推理
bool RMSNormRel(const Array<Type>& types, int num_inputs,const Attrs& attrs, const TypeReporter& reporter){
  ICHECK_EQ(types.size(),3); // [data,gamma,output] 数量为输入+输出
  const auto* data=types[0].as<TensorTypeNode>();
  if(data==nullptr) return false;
  const RMSNormAttrs* param = attrs.as<RMSNormAttrs>();
  int axis = param->axis>=0 ? param->axis: param->axis+data->shape.size(); //axis可能是python风格的负数 如-1
  ICHECK(axis >= 0 && axis < (int)data->shape.size());
  reporter->Assign(types[1], TensorType({data->shape[axis]}, data->dtype)); //gamma的shape与axis所在shape相同
  reporter->Assign(types[2], TensorType({data->shape}, data->dtype));
//output的shape与data的shape相同
  return true;
}

算子与属性关联


注册算子的名称及其他描述,并用调用接口进行标注。

在tvm/src/relay/op/nn/http://nn.cc添加以下代码:
//根据输入与属性调用算子
Expr MakeRMSNorm(Expr data, Expr gamma,int axis,double epsilon,bool scale) {
  auto attrs=make_object<RMSNormAttrs>();
  attrs->axis=axis;
  attrs->epsilon=epsilon;
  attrs->scale=scale;
  static const Op& op = Op::Get("nn.RMS_norm");
  return Call(op,{data,gamma}, Attrs(attrs), {});
}

TVM_REGISTER_GLOBAL("relay.op.nn._make.RMS_norm").set_body_typed(MakeRMSNorm);

//注册算子
RELAY_REGISTER_OP("nn.RMS_norm")
  .describe(R"code(
    RMSNorm: It is a replacement of LayerNorm.
        Zhang B, Sennrich R. Root mean square layer normalization[J]. Advances in Neural Information Processing Systems, 2019, 32
    )code" TVM_ADD_FILELINE)
    .set_attrs_type<RMSNormAttrs>()
    .set_num_inputs(2)
    .add_argument("data","Tensor","Input to which RMS_norm will be applied")
    .add_argument("gamma","Tensor","The gamma scale factor.")
    .set_attr<FInferCorrectLayout>("FInferCorrectLayout",  NormalizationInferCorrectLayout<RMSNormAttrs>)
    .set_support_level(1)
    .add_type_rel("RMSNorm",RMSNormRel);

定义算子的计算


TVM的TOPI算子库包含多个后端的算子的计算与调度定义,在这里进行算子在Python端注册。

新增tvm/python/tvm/topi/nn/RMS_norm.py文件,编写如下代码:

利用TVM的跨语言调用机制,将RMSNorm的计算定义编写在CPP端,在Python端提供接口
"""RMS normalization operator."""
from .. import cpp

def RMS_norm(data, gamma, axis, epsilon=1e-5):
    """RMS normalization operator.

    Parameters
    ----------
    data : tvm.te.Tensor
        N-D with shape (d_0, d_1, ..., d_{N-1})

    gamma: tvm.te.Tensor
        K-D with shape (r_0, r_1, ..., r_{K-1}) where K == len(axis) and d_{axis_k} == r_k

    axis : list of int
        Axis over the normalization applied

    epsilon : float
        The epsilon value to avoid division by zero.

    Returns
    -------
    result : tvm.te.Tensor
        N-D with shape (d_0, d_1, ..., d_{N-1})
    """
    return cpp.nn.RMS_norm(data, gamma, axis, epsilon)

注意在tvm/python/tvm/topi/nn/_ init _.py导入算子

新增tvm/include/tvm/topi/nn/RMS_norm.h文件,编写如下代码:

该代码描述了RMSNorm的计算流程
/*!
 * \brief RMS normalization op constructions
 * \file nn/RMS_norm.h
 */
#ifndef TVM_TOPI_NN_RMS_NORM_H_
#define TVM_TOPI_NN_RMS_NORM_H_

#include <tvm/te/operation.h>
#include <tvm/topi/tags.h>

#include <string>

namespace tvm {
namespace topi {
namespace nn {

using namespace tvm::te;

/*!
 * \brief RMS normalization.
 * \param data N-D tensor with shape [d_0, d_1, ..., d_{N-1}]
 * \param gamma K-D tensor with shape [r_0, r_1, ..., r_{K-1}] where K == len(axis) and
 *              d_{axis_k} == r_k
 * \param axis The axis to normalize over.
 * \param epsilon The epsilon value to avoid division by zero.
 * \param name The name of the operation.
 * \param tag The tag to mark the operation.
 * \return The normalized tensor, with the same shape as data.
 */
inline Tensor RMS_norm(const Tensor& data, const Tensor& gamma,
                         const Array<Integer>& axis, double epsilon,
                         std::string name = "T_RMS_norm", std::string tag = kInjective) {
  // sum  x^2
  auto ndim = data->shape.size();
  ICHECK_NE(ndim, 0) << "Cannot reduce a 0 dim Tensor";
  auto real_axis = GetRealAxis(static_cast<int>(ndim), axis);
  auto reduce_axes = MakeReduceAxes(real_axis, data);
  auto target_shape =
      MakeReduceTargetShape(real_axis, data, /*keepdims=*/false, /*atleast1d=*/true);
  auto func = MakeTupleSumReducer();

  auto compute = [ndim, &real_axis, &reduce_axes, &func, &data](const Array<Var>& indices) {
    Array<PrimExpr> eval_range;
    int arg_counter = 0;
    int red_counter = 0;

    for (size_t i = 0; i < ndim; ++i) {
      if (std::find(real_axis.begin(), real_axis.end(), i) != real_axis.end()) {
        // real_axis contains i
        eval_range.push_back(reduce_axes[red_counter]);
        red_counter++;
      } else {
        eval_range.push_back(indices[arg_counter]);
        arg_counter++;
      }
    }
    auto square = [](const PrimExpr& x) { return x * x; };
    return func({data(eval_range), square(data(eval_range))}, reduce_axes, nullptr);
  };

  auto temp_x_x2 =
      tvm::te::compute(target_shape, compute, data->op->name + "_red_temp", kCommReduce);

//获得平方和
  auto temp_x2 = temp_x_x2[1];

//平方和求均值时要除以的元素的数量
  auto reduce_extent = make_const(data->dtype, 1);
  for (int i : real_axis) {
    reduce_extent *= data->shape[i];
  }

  auto RMS_norm_func = [&](const Array<Var>& indices) {
    Array<Var> reduce_indices, non_reduce_indices;
    for (int i = 0, n = static_cast<int>(indices.size()); i < n; ++i) {
      if (std::find(real_axis.begin(), real_axis.end(), i) != real_axis.end()) {
        reduce_indices.push_back(indices[i]);
      } else {
        non_reduce_indices.push_back(indices[i]);
      }
    }

    auto var = temp_x2(non_reduce_indices) / reduce_extent ;
    auto RMS_norm = data(indices) * tvm::rsqrt(var + make_const(var->dtype, epsilon));  //tvm::rsqrt 即 1/tvm::sqrt
    RMS_norm = topi::multiply(RMS_norm, gamma(reduce_indices));

    return RMS_norm;
  };

  return tvm::te::compute(data->shape,RMS_norm_func, name, tag);
}

}  // namespace nn
}  // namespace topi
}  // namespace tvm

#endif  // TVM_TOPI_NN_RMS_NORM_H_

在tvm/src/topi/http://nn.cc增加如下代码:

注册topi算子库的CPP端
// 注意引入下面的头文件
#include <tvm/topi/nn/RMS_norm.h>

/* Ops from nn/RMS_norm.h */
TVM_REGISTER_GLOBAL("topi.nn.RMS_norm").set_body([](TVMArgs args, TVMRetValue* rv) {
*rv = nn::layer_norm(args[0], args[1], args[2], static_cast<double>(args[3]));
});

提供Python API


在tvm/python/tvm/relay/op/nn/nn.py增加如下代码:
def RMS_norm(data,gamma,axis=-1,epsilon=1e-5,scale=True):
    return _make.RMS_norm(data,gamma,axis,epsilon,scale)

编写测试文件


新增tvm/python/tvm/topi/testing/RMS_norm_python.py,编写如下代码:

该文件是用于topi算子的测试,使用numpy编写,在测试时作为标准答案
import numpy as np
from functools import reduce

def RMS_norm_python(data, gamma, axis, epsilon=1e-5):
    """RMS normalization operator in Python.

    Parameters
    ----------
    data : numpy.ndarray
        N-D with shape (d_0, d_1, ..., d_{N-1})

    gamma: numpy.ndarray
        K-D with shape (r_0, r_1, ..., r_{K-1}) where K == len(axis) and d_{axis_k} == r_k

    axis : int or tuple of ints
        Axis over the normalization applied

    epsilon : float
        The epsilon value to avoid division by zero.

    Returns
    -------
    result : np.ndarray
        N-D with shape (d_0, d_1, ..., d_{N-1})
    """

    if len(axis)==1:
        n=data.shape[axis[0]]
    else:
        n=reduce(lambda x,y:data.shape[x]*data.shape[y],axis)

    temp=np.sum(np.square(data),axis) / n
    temp=np.repeat(temp,axis=0,repeats=n).reshape(data.shape)
    result = data/ np.sqrt(temp + epsilon)
    result *= gamma

    return result

注意在tvm/python/tvm/topi/testing/init.py导入,导入代码如下
from .RMS_norm_python import RMS_norm_python

新增tvm/tests/python/topi/python/test_topi_RMS_norm.py,编写如下代码
"""Test code for RMS_norm."""
import numpy as np
import pytest
import tvm
from tvm import te
from tvm import topi
from tvm.topi.utils import get_const_tuple
import tvm.topi.testing

import tvm.testing

# 使用通用的injective调度
_RMS_norm_schedule = {
    "generic": topi.generic.schedule_injective,
}

# 对最后一维和最后两维分别测试
@tvm.testing.parametrize_targets("llvm")
@pytest.mark.parametrize("shape,axis", [([4, 16], (1,)), ([4, 16, 16], (1, 2))])
def test_layer_norm(target, dev, shape, axis, episilon=1e-5, dtype="float32", rtol=1e-5, atol=1e-5):
    data = te.placeholder(shape, dtype=dtype, name="data")
    scale_shape = [shape[dim] for dim in axis]
    gamma = te.placeholder(scale_shape, dtype=dtype, name="gamma")
    B = topi.nn.RMS_norm(data, gamma, axis, episilon)# 调用TOPI算子库中的RMSNorm

    data_np = np.random.uniform(size=shape).astype(dtype)
    gamma_np = np.random.uniform(size=scale_shape).astype(dtype)
    beta_np = np.random.uniform(size=scale_shape).astype(dtype)
    b_np = tvm.topi.testing.RMS_norm_python(data_np, gamma_np, axis, episilon) # 调用numpy编写的RMSNorm作为标准答案

    with tvm.target.Target(target):
        s_func = tvm.topi.testing.dispatch(target, _RMS_norm_schedule)
        s = s_func([B])
    data_tvm = tvm.nd.array(data_np, dev)
    gamma_tvm = tvm.nd.array(gamma_np, dev)
    b_tvm = tvm.nd.array(np.zeros(get_const_tuple(B.shape), dtype=dtype), dev)
    f = tvm.build(s, [data, gamma, B], target)
    f(data_tvm, gamma_tvm, b_tvm)
    tvm.testing.assert_allclose(b_tvm.asnumpy(), b_np, rtol=rtol, atol=atol)

if __name__ == "__main__":
    tvm.testing.main()

执行测试


由于涉及到CPP代码修改,因此需要重新make整个项目生成更新的动态链接库。

之后执行tvm/tests/python/topi/python/test_topi_RMS_norm.py测试文件,得到以下结果:表示测试成功

AI深度学习编译器工程师需要哪些技术栈?6854 作者: 来源: 发布时间:2023-8-5 00:16

总结


TVM中算子的注册并不局限这一种方法,例如对于BatchNorm,并没有在CPP端描述其计算,而是在python端定义(tvm/topi/nn/batch_norm.py),对于其调度,也有不同后端的多个实现。

本文只包含了编译器中新增算子,后续将更新在Pytorch中定义RMSNorm、导出为ONNX、利用Relay前端导入的流程。(挖坑:)


AI深度学习编译器工程师需要哪些技术栈?5706 作者: 来源: 发布时间:2023-8-5 00:16
https://www.zhihu.com/question/532768471/answer/2682461714

相关从业者,持续补充基础知识,共勉:
  • mlir/llvm(很多招聘职位上都明确写着有LLVM、MLIR开发经验者优先考虑~)
  • pytorch/tensorflow等一个或多个深度学习框架,会自己构造一些小模型跑测试
  • onnx
  • 语言:python,C++(C++11,模板)
  • 数学:线性代数(矩阵乘法,卷积的基础)
  • 深度学习基础知识:卷积、池化、激活、前向传播反向传播等基础操作,Resnet、MobileNet、Inception、vgg、yolo等经典网络系列等
  • 编译原理,IR/Pass(编译器的优化是无止境的,到后期会真正体会到)
  • 操作系统原理:访存、存储层次结构
  • 计算机体系结构:指令集(有的算子发射可能涉及到)、流水线
  • 其他:git管理/linux常用命令
  • 如果是做CodeGen(即算子发射)这一块,则需要了解NPU(DPU、GPU、_PU)或者SoC硬件模块单元功能、寄存器配置、DDR/SRAM访存特性等硬件特性、运算流程和限制等知识。
  • 如果是做前端,则需要了解xla、onnx、mlir等

推荐书籍:《智能计算系统》
点击👇卡片关注我,第一时间获取干货~



上一篇:关于四川会计人员提升学历的重要通知!
下一篇:宝塔面板-未授权访问任意登入

鲜花

握手

雷人

路过

鸡蛋

评论

您需要登录后才可以发表言论 登录立即注册
创宇盾启航版免费网站防御网站加速服务
投诉/建议联系

discuzaddons@vip.qq.com

未经授权禁止转载,复制和建立镜像,
如有违反,按照公告处理!!!
  • 联系QQ客服
  • 添加微信客服

联系DZ插件网微信客服|最近更新|Archiver|手机版|小黑屋|DZ插件网! ( 鄂ICP备20010621号-1 )|网站地图 知道创宇云防御

您的IP:3.144.105.101,GMT+8, 2024-12-23 19:14 , Processed in 0.092453 second(s), 43 queries , Gzip On, Redis On.

Powered by Discuz! X5.0 Licensed

© 2001-2024 Discuz! Team.

关灯
扫一扫添加微信客服
QQ客服返回顶部
返回顶部