精华内容
下载资源
问答
  • 1. TVM编译TVM4J-tvm4j-core生成的本地MAVEN库 2. 解压到本地用户目录(.m2目录),然后使用AndroidStudio编译TVM的app目录下的三个Android APK(android-camera, android-deploy,android-rpc) 3. 需要在Linux环境编译...
  • 深度学习编译引擎TVM. 也可以加群讨论: 61719120x (不有广告,最后一个用x代表,有心的人自己猜群号吧) TVM 0.8版起新api测试代码,模型文件 all in one, 编译好后就可以测试使用。 使用全新API. 无需param/json,...
  • g++ -std=c++11 -g t_m.cpp -o out -L /usr/local/lib -lopencv_core -lopencv_imgproc -lopencv_highgui -lopencv_imgcodecs -lopencv_dnn -L /home/u260260/mntspace/dowloads/apache-tvm-src-v0.6.1.rc1-...
  • Neo-AI / TVMTVM的下游分支,在上游代码库的顶部包括特定于供应商和产品的功能。 分行 开发者 -这是具有最新源代码更新的开发分支。 执照 :copyright:贡献者根据许可获得许可。 贡献TVM TVM采用apache提交者模型...
  • 更新(20200429) 此存储库基于库。 将这个高效的付款网络转换为ONNX的需求很多,因此我们进行此回购以帮助人们将模型转换为ONNX或TVM。 请注意,此存储库仅提供如何将模型转换为... 然后,在tvm_home/python/tvm/rela
  • 3DNN-FPGA-TVM-源码

    2021-02-26 06:03:30
    3DNN-FPGA-TVM #使用TVM编译器在FPGA板上实现3D计算机视觉深度学习
  • 陈天奇给了关于TVM的报告,TVM: An End to End Automated Deep Learning Compiler。
  • TVM编译框架

    2018-12-25 21:21:25
    TVM 编译框架的论文,
  • TVM和VTA的集体知识存储库 所有CK组件都可以在和! 该项目由托管。 介绍 该存储库为提供了高级,可移植且可自定义的集体知识工作流。 这是我们长期社区计划的一部分,该计划使用统一和自动化AI,ML和系统R&D,并在...
  • TVM是深度学习系统的编译器堆栈。 它旨在缩小以生产力为重点的深度学习框架与以性能和效率为重点的硬件后端之间的差距。 TVM与深度学习框架协同工作,为不同的后端提供端到端编译。
  • TVM链接器 该存储库存储tvm_linker实用程序的源代码。 它需要TON智能合约的TVM( )汇编源代码,对其进行编译和链接,添加标准选择器和运行时,并将其存储到二进制TVC文件中。 而且,它可以通过模拟TON交易的计算...
  • TVM计算器 曾经不得不上过工程经济学课程吗? 直到我在UBC修读必修课程,我才知道它甚至不存在。 我发现总是手动计算比率很麻烦,而是创建了一个api(待完成的Web和Android应用程序)来替您完成计算! 介绍.... ...
  • 用于TVM的C和C ++编译器用于TVM的C和C ++编译器基于LLVM 7.0.0的TVM 0.2.0的Clang。 该存储库存储用于TON虚拟机的Clang源代码和运行时库,该库随后将移至单独的存储库。 TON VM(https://test.ton.org/tvm.pdf)旨在...
  • 金融计算器(tvm、npv、现金流、简单计算器、房贷摊销计算器等,利率转换等8中计算器
  • 文件包括:操作软件,xbox360摄像头驱动,视频教程,pdf说明书。
  • 上一篇《TVM的“hello world“基础流程 I》中基于一个最基本的case介绍了TVM中计算的定义与schedule的构建。这篇沿用上一篇中的case,继续介绍接下去的一个重点部分,就是编译。 编译 有了前面构建的schedule之后,...

    上一篇《TVM的“hello world“基础流程 I》中基于一个最基本的case介绍了TVM中计算的定义与schedule的构建。这篇沿用上一篇中的case,继续介绍接下去的一个重点部分,就是编译。

    有了前面构建的schedule之后,接着就需要编译并生成目标代码了。这个工作主要由tvm.build()relay.build()两个函数来完成。它们的区别在于应用目标的范围,前者用于单个算子,后者用于整个网络。由于网络可看作由算子组成,后者会调用前者。本例中是针对单个算子的,因此这里使用的是前者:

    tgt = tvm.target.Target(target="llvm", host="llvm")
    fadd = tvm.build(s, [A, B, C], tgt, name="vecadd")
    

    其中最主要的build()函数定义在driver/build_module.py文件中。该函数基于给定参数构建出可调用的目标函数。按照官方介绍里的说法,它主要做两个工作 :

    • Lowering:将high-level的循环嵌套结构转换成最终的low-level的IR。
    • Codegen:从low-level的IR生成目标机器代码。

    该函数的第一个参数是前面构建出来的schedule,第二个参数是函数的参数列表,第三个参数是target。它提供用于lowering和codegen所需的目标平台信息。代码中对应的Target对象定义在target.*文件中。其构造函数有两个参数,其中第一个参数target指示目标平台的配置。其中的配置项比如:

    • kind: 平台类型,它基本决定了生成的代码是在什么处理器上运行。注册的target kind详细见target_kind.cc,有llvm, c, cuda, nvptx, romc, opencl, metal, vulkan, hexagon等。
    • keys: 如kind是opencl的话,key可以是mali, opencl, gpu。
    • device:对应实际运行的设备,它会添加到keys后面。
    • libs:外部库,如cblas, cudnn, cublas, mkl这些。

    另外参数host与target类似,但它用于指示host平台。比如taret平台为cuda的话,毕竟GPU还是不能完全脱离CPU运行,因此还需要host的代码做胶水,如内存分配,kernel启动这些。默认为llvm。

    Lowering过程可以单独用tvm.lower()函数完成,如:

    m = tvm.lower(s, [A, B, C], name="vecadd")
    rt_mod = tvm.build(m, target="llvm")
    

    也可以通过tvm.build()函数完成(因为它一进去就会先调用lower()函数)。lower()函数的主要流程相关代码:

    lower(sch, args, name="main", ...) // driver/build_module.py
        // Handle add_lower_pass, if any.
        lower_phases0 = ...
        ...
        // According to the given schedule, form a function (in IRModule).
        mod = form_irmodule(sch, args, ...) // build_module.py
            sch.normalize()
                Schedule::normalize() // schedule_dataflow_rewrite.cc
                    InjectInline()
                    RebaseNonZeroMinLoop()
                    LegalizeInvalidAttach()
            bounds = schedule.InferBound(sch)  
                InferBound() // bound.cc
            stmt = schedule.ScheduleOps(sch, bounds)
                ScheduleOps() // schedule_ops.cc
                    body = Stmt()
                    // scan init and scan updates
                    ...
                    for each stage in schedule: // in reverse order
                        body = MakePipeline(stage, dom_map, body, ...)
                    SchedulePostProc post_proc
                    post_proc.Init(sch)
                    return post_proc(body)
            compact = schedule.VerifyCompactBuffer(stmt)
            binds, arg_list = get_binds(args, compact, binds)
            stmt = schedule.SchedulePostProcRewriteForTensorCore(stmt, sch, ...)
            // func type: PrimFunc
            func = schedule.SchedulePostProcToPrimFunc(arg_list, stmt, ...) // schedule_postproc_to_primfunc.cc
                // Prepare parameters
                ...
                return tie::PrimFunc(params, body, ...)
            // name: vecadd
            func = func.with_attr("global_symbol", name)
            // Set functions
            return tvm.IRModule({name: func})
            
        // Phase 0: InjectPrefetch, StorageFlatten, BF16Legalize, NarrowDataType, Simplify
        pass_list = lower_phase0
        
        // Phase 1: LoopPartition, VectorizeLoop, InjectVirtualThread, InjectDoubleBuffer, StorageRewrite, UnrollLoop
        pass_list += lower_phase1
        
        // Phase 3: Simplify, RemoveNoOp, RewriteUnsafeSelect, HoistIfThenElse
        pass_list += lower_phase2
        
        // Apply the above passes.
        optimize = tvm.transform.Sequential(pass_list)
        mod = optimize(mod) 
        
        // mod type: tvm.ir.module.IRModule
        return mod 
    

    它主要根据参数给的schedule与参数生成对应的IRModule对象(定义在ir/module.h中)。IRModule是软件栈中所有IR变换的基础单元。它维护函数与类型定义。这里的各种pass就是在IRModule上进行并吐出IRModule

    在这里插入图片描述

    其中几个主要数据结构关系如下:
    在这里插入图片描述

    lower()函数中有四个阶段,第一个阶段中通过form_irmodule()函数根据给定的schedule生成IRModule对象,然后在这个IRModule对象上应用4轮的pass。这些pass主要分为几个阶段,分别是:

    • Phase 0:使用者自定义的pass。
    • Phase 1:使用者自定义的pass。以及:
      • InjectPrefetch
      • StorageFlatten
      • BF16Legalize
      • NarrowDataType
      • Simplify
    • Phase 2:使用者自定义的pass。以及:
      • LoopPartition
      • VectorizeLoop
      • InjectVirtualThread
      • InjectDoubleBuffer
      • StorageRewrite
      • UnrollLoop
    • Phase 3:使用者自定义的pass。以及:
      • Simplify
      • RemoveNoOp
      • RewriteUnsafeSelect
      • HoistIfThenElse
      • InstrumentBoundCheckers

    这此pass其实是编译构建过程中的精华之一。但限于篇幅(其实是我自己也没了解全。。。),以后再进一步讨论。

    lower()函数的最后返回经过上面多轮pass优化后的IRModule对象。其中form_irmodule()函数是相对比较复杂的一部分,它主要负责生成最初的IRModule对象,其中几个关键步骤如下:

    1. Schedule::normalize()函数规范化给定的schedule。主要实现在schedule_dataflow_rewrite.cc文件中。它调用以下三个函数。本例比较简单,因此它们实际都没有起什么作用。。。
      1. InjectInline()函数处理算子内联。用到调度原语 compute_inline的话会用到。
      2. RebaseNonZeroMinLoop()函数将循环迭代的最小界置为0。感觉有点canonicalization的意思。
      3. LegalizeInvalidAttach()函数处理在使用调度原语compute_at时且目标迭代又被split或fuse情况下的合法化。
    2. InferBound()函数顾名思义就是边界推导(Bound inference),主要用于推导循环边界。更具体地,就是确定每个IterVar的范围,它返回IterVarRange的映射,即每个循环变量的范围。这个信息在后面的MakeLoopNest()函数中用于确定for循环的范围,和在BuildRealize()函数中设置缓冲的大小。具体可参见官方文档 InferBound Pass
    3. ScheduleOps()函数基于前面经过一些处理后的Schedule对象和推导出来的循环边界产生Stmt对象。它表示一个初始的循环嵌套结构。C++层中的Stmt为所有语句(Statement)的容器。它的子类有LetStmtAttrStmtAssertStmtStoreAllocateSeqStmtIfThenElseEvaluateForWhile等等。该函数会处理schedule的依赖,核心部分是逆向遍历Schedule当中的Stage(对于上面例子中就是先Compute Op,再两个Placeholder Op)。对于每个stage(PlaceholderOp除外),根据其attach type调用相应的逻辑。
      1. 对于上面的例子,Compute Op没有attach在其它计算中,因此它对应Stage的attach type为kGroupRoot,因此这里调用MakePipeline()函数产生Stmt。这步比较关键比较复杂,后面再展开。
      2. 然后通过SchedulePostProc对象(继承自StmtExprMutator)对前面生成的Stmt进行后处理。
    4. get_binds()函数用于绑定buffer。它给每个参数张量分配buffer。如对于上面例子中的A, B, C三个张量,分别通过tvm.tir.decl_buffer()创建buffer并将之与张量绑定。
    5. SchedulePostProcToPrimFunc()函数基于ScheduleOps()产生的Stmt创建PrimFunc对象,它可以被用于TIR优化。PrimFunc代表包含了TIR statement的primitive function,它是low-level的代码表示。
    6. 创建IRModule对象。基于上面生成的对象封装成IRModule对象并返回。一个IRModule可以有多个函数,比较简单的情况下就一个。

    上面第ScheduleOps()函数中会调用MakePipeline()函数针对ComputeOp对应Stage,返回一条由Stmt组成的pipeline,其大体流程相关代码如下:

    MakePipeline(Stage, unordered_map<IterVar, Range>, Stmt, ...) // schedule_ops.cc
        producer = s->op->BuildProvide(stage, ...) // ComputeOpNode::BuildProvide() in compute_op.cc
            ComputeType ctype = DetectComputeType(this, stage)
            MakeComputeStmt(...) // compute_op.cc
                ComputeLoopNest n = ComputeLoopNest::Create(...) // compute_op.cc
                    ComputeLoopNest ret
                    // make main loop nest
                    ret.main_nest = MakeLoopNest(stage, dom_map, ...) // op_utils.cc
                        vector<vector<Stmt>> nest
                        nest.resize(leaf_iter_vars.size() + 1)
                        for iter_var in leaf_iter_vars:
                            nest[i + 1].emplace_back(For(var, 0, dom->extent, kind, no_op))
                            nest[i + 1].emplace_back(AttrStmt(iv, tir::attr::loop_scope, iv->var, no_op))
                    ...
                n.init_nest.emplace_back(MakeIfNest(n.init_predicates))
                n.main_nest.emplace_back(MakeIfNest(n.main_predicates))
                if has reduce_axis:
                    ...
                else:
                    vector<Stmt> provides
                    ...
                    // Array<Stmt> -> SeqStmt
                    Stmt provide = SeqStmt::Flatten(provides) // stmt.h
                    provide = MergeNest(n.main_nest, provide) // ir_utils.cc
                    return Substitute(provide, n.main_vmap) // stmt_functor.cc
        Stmt pipeline = producer
        pipeline = s->op->BuildRealize(stage, dom_map, pipeline) 
            // set the sizes of allocated buffers
            BaseComputeOpNode::BuildRealize(stage, realize_map, body) // compute_op.cc
                Stmt realize = body
                realize = tir::ProducerRealize(...)
        pipeline = AttrStmt(s->op, tir::attr::realize_scope, ..., pipeline)
        return pipeline
    

    MakePipeline()函数主要步骤如下:

    1. ComputeOpNode::BuildProvide()函数主要创建ComputeOp对应的循环嵌套对应的那些Stmt对象并串成pipeline。
      1. 首先用DetectComputeType()函数检测计算类型。它遍历当前Stage的所有当前有效IterVar对象,并根据它们的属性判断计算类型,对于上面的简单例子这里为ComputeType::kNormal
      2. 然后根据类型调用相应函数创建Stmt对象。这里对应地是调用MakeComputeStmt()函数。
        1. 根据Stage对象和边界推导的结果通过ComputeLoopNest::Create()函数创建ComputeLoopNest对象。该对象表示循环嵌套,它几个主要成员:

          • init_predicatesmain_predicates:类型为vector<PrimExpr>。表示每个循环的边界判断,调用MakeBoundCheck()函数来生成。
          • init_nestmain_nest:类型为vector<vector<Stmt>>。 其中main_nest是最主要的表示循环嵌套的对象,对于上面的例子,经过split后这里包含两个for循环。
        2. 根据main_predicates创建对应的Stmt(如有),用于在循环中判断该predicate是否成立,并添加到main_nest结构中。

        3. 根据有无reduce axis走不同的path。如果没有的话(如本例),对于ComputeOpbody中的每一个输出,创建ProducerStore对象,再通过MergeNest()函数将之与主嵌套main_nest合并。

        4. 通过Substitute()函数基于main_vmap(在MakeLoopNest()函数中准备)进行替换。

    2. 如schedule中设置了double buffer(如s[A].double_buffer),则添加对应的AttrStmt。它通过增大额外的buffer来达到达到计算与访存的重叠。本例中没用到。
    3. 如传入的consumer有定义且不是no op(指无定义、const init的EvaluateNode,或者是长度为0的SeqStmtNode),则添加SeqStmtproducerconsumer串连起来。本例中也不适用。
    4. 调用BuildRealize()函数。对于每个输出的张量,在pipeline中加入ProducerRealize节点。
    5. 最后,在pipeline中添加AttrStmt节点标注操作的范围,并返回该pipeline。

    对于前面vecadd的例子,得到的pipeline大致如下示意图:
    在这里插入图片描述

    整个lower()函数后完成后的IR(TIR)打印出来如下:

    primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
      attr = {"global_symbol": "main", "tir.noalias": True}
      buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024], []),
                 B: Buffer(B_2: Pointer(float32), float32, [1024], []),
                 A: Buffer(A_2: Pointer(float32), float32, [1024], [])}
      buffer_map = {A_1: A, B_1: B, C_1: C} {
      for (i.outer: int32, 0, 16) {
        for (i.inner: int32, 0, 64) {
          C_2[((i.outer*64) + i.inner)] = ((float32*)A_2[((i.outer*64) + i.inner)] + (float32*)B_2[((i.outer*64) + i.inner)])
        }
      }
    }
    

    Lowering完成后,接下去就是build了。Build的主要流程相关代码如下:

    build() # driver/build_module.py
        input_mod = lower(inputs, args, ...) 
    
        mod_host_all = tvm.IRModule()
    
        for tar, input_mod in target_input_mod.items():
            # build the lowered functions for a device with the given compilation
            mod_host, mdev = _build_for_device(input_mod, tar, target_host)
                # input_mod type: IRModule
                mod_mixed = input_mod 
                # Apply passes:  ThreadSync, InferFragment, LowerThreadAllreduce, MakePackedAPI, SplitHostDevice
                ...
                # Device optimizations: Filter, LowerWarpMemory, ,Simplify, LowerDeviceStorageAccessInfo, LowerIntrin
                ...
                mod_dev = opt_device(mod_mixed) # IRModule
                # Host optimization: LowerTVMBuiltin, LowerDeviceStorageAccessInfo, CustomDataType, LowerIntrin, CombineContextCall
                ...
                mod_host = opt_host(mod_mixed) # IRModule
                
                # Build IRModule into Module
                # If there are dev functions
                rt_mod_dev = codegen.build_module(mod_dev, target) # target/codegen.py
                    _ffi_api.Build(mod, target) # codegen.py
                # mod_host type: IRModule, rt_mod_dev type: Module
                return mod_host, rt_mod_dev 
            mod_host_all.update(mod_host)
                # Insert functions in another Module to current one
                _ffi_api.Module_Update()
                    IRModuleNode::Update() # ir/module.cc
            device_modules.append(mdev)
        # Generate a unified host module (type: runtime.Module)
        rt_mod_host = codegen.build_module(mod_host_all, target_host)
            # Create LLVMModuleNode and return the corresponding Module
            _ffi_api.Build(mod, target) # target/codegen.cc
        # Import all modules
        for mdev in device_modules:
            rt_mod_host.import_module(mdev)
                _LIB.TVMModImport(mod, dep) # c_runtime_api.cc
                    GetModuleNode(mod)->Import(...) # runtime/module.cc
                        imports_.emplace_back(...)
        return rt_mod_host # runtime.module.Module
    

    target_input_mod包含了前面lowering输出的需要编译的IRModule及相应的target信息。比如LLVM(CPU)为target,就是:{"llvm -keys=cpu -link-params=0", IRModule}。如cuda为target,可能就是{“cuda -keys=cuda,gpu -max_num_threads=1024 -thread_warp_size=32", IRModule}。对于简单的case(如本文这个),target_input_mod只包含一个元素,_build_for_device()函数返回host端的IRModule,以及target端的Module(如是cuda平台的话C++层对应CUDAModuleNode对象)。然后将host端IRModule生成一个统一的host模块,再将前面生成的对应target的Module(如有)导入其中。

    这里,其中mod_host_allmod_host的类型为tvm.ir.module.IRModulert_mod_hostmdev的类型为tvm.runtime.module.Module。注意mdev只有当目标为非CPU(如GPU等)平台时才会有,当target为llvm(即for CPU)时mdev为空。

    这个流程大体示意图如下:
    在这里插入图片描述

    其中比较核心和重要的部分是Build()函数,实现在codegen.cc文件中。它会调用到具体后端的编译函数,进行目标代码生成。如cuda平台的话对应函数定义在build_cuda_on.cc文件中,llvm的话在llvm_module.cc文件中。以llvm后端为例,其主要流程相关代码为:

    TVM_REGISTER_GLOBAL("target.build.llvm")
        .set_body_typed([](IRModule mod, Target target) -> runtime::Module { 
            auto n = make_object<LLVMModuleNode>();
            n->Init(mod, target); // llvm_module.cc
                InitializeLLVM();
                    llvm::InitializeAllTargetInfos();
                    llvm::InitializeAllTargets();
                    ...
                unique_ptr<CodeGenLLVM> cg = CodeGenLLVM::Create(...) // codegen_llvm.cc
                    // Call the corresponding codegen backend according to the target.
                    const PackedFunc* f = runtime::Registry::Get("tvm.codegen.llvm.target_" + target);
                    handle = (*f)() 
                    return unique_ptr<CodeGenLLVM>(handle);
                    
                vector<PrimFunc> funcs;
                for kv : mod->functions:
                    ...
                    f = Downcast<PrimFunc>(kv.second);
                    if (f->HasNonzeroAttr(tir::attr::kIsEntryFunc))
                        entry_func = global_symbol.value();
                    funcs.push_back(f);
                cg->Init("TVMMod", ...);
                    CodeGenCPU::Init() // codegen_cpu.cc
                        CodeGenLLVM::Init() // codegen_llvm.cc
                        
                for f in funcs:
                    cg->AddFunction(f); // codegen_cpu.cc
                        CodeGenLLVM::AddFunction();
                            AddFunctionInternal(f);
                                llvm::FunctionType* ftype = llvm::FunctionType::get(...);
                                // kGlobalSymbol: "global_symbol"
                                global_symbol = f->GetAttr<String>(tvm::attr::kGlobalSymbol);
                                function_ = llvm::Function::Create(...);
                                llvm::BasicBlock* entry = llvm::BasicBlock::Create(..., function_);
                                IRBuilder::SetInsertPoint(entry);
                                this->VisitStmt(f->body);
                                builder_->CreateRet(ConstInt32(0));
                if entry_func.length() != 0:
                    cg->AddMainFunction(entry_func); // codegen_cpu.cc
                        // tvm_module_main : "__tvm_main__"
                        llvm::GlobalVariable* global = new llvm::GlobalVariable(*module_, ..., tvm_module_main);
                        global->setInitializer(llvm::ConstantDataArray::getString(*ctx_, entry_func_name))
                        global->setDLLStorageClass(llvm::GlobalVariable::DLLExportStorageClass);
                module_ = cg->Finish(); // CodeGenCPU::Finish() in codegen_cpu.cc
                    CodeGenLLVM::Finish(); // codegen_llvm.cc
                        CodeGenCPU::AddStartupFunction();
                            function_ = llvm::Function::Create(ftype, llvm::Function::InternalLinkage,"__tvm_module_startup", module_.get());
                            llvm::BasicBlock* startup_entry = llvm::BasicBlock::Create(*ctx_, "entry", function_);
                            llvm::appendToGlobalCtors(*module_, function_, 65535);
                            builder_->CreateRet(nullptr);
                        CodeGenLLVM::Optimize(); // codegen_llvm.cc
                            // Function pass manager
                            FPassManager fpass(module_.get());
                            // Module pass manager
                            MPassManager mpass;
                            mpass.add(llvm::createTargetTransformInfoWrapperPass(getTargetIRAnalysis()));
                            fpass.add(llvm::createTargetTransformInfoWrapperPass(getTargetIRAnalysis()));
                            llvm::PassManagerBuilder builder;
                            builder.Inliner = llvm::createFunctionInliningPass(builder.OptLevel, ...);
                            builder.LoopVectorize = true; 
                            builder.SLPVectorize = true; 
                            ...
                            // Run the function passes
                            for mod in module_:
                                fpass.run(mod);
                            fpass.doFinalization();
                            // Run the module passes.
                            mpass.run(*module_);
            return runtime::Module(n);
        });
    

    该函数中先创建LLVMModuleNode对象,然后调用它的Init()函数进行初始化,最后封装成Module对象返回。其中的Init()函数主要是将之前生成的TIR转为LLVM IR。它主要分几步:

    1. InitializeLLVM()函数初始化LLVM环境。这里边主要是例行调用LLVM的一大坨初始化函数。

    2. 创建用于代码生成的CodeGenLLVM对象。这里由于target字符串为x86-64,因此工厂函数名为tvm.codegen.llvm.target_x86-64。该工厂函数中创建CodeGenX86_64对象。因为继承关系为CodeGenX86_64 -> CodeGenCPU -> CodeGenLLVM,所以返回的是CodeGenLLVM的指针。

    3. 类型为IRModule的参数mod中的functions成员包含了该模块中的函数。这一步中将这些函数存于类型PrimFunc的数组funcs中。对于标为入口函数(kIsEntryFunc)的函数,记录在entry_func变量中。

    4. 接下来初始化前面创建的CodeGenX86_64对象。先调用CodeGenCPU::Init(),它里边又会调用到CodeGenLLVM::Init()。前者主要创建一坨TVM运行时类型与函数。后者创建一些llvm中用于codegen的对象,如IRBuilderllvm::Modulellvm::MDBuilder

    5. 对前面放入funcs数组的每个函数,调用CodeGenCPU::AddFunction()函数进行代码生成。对本文涉及的case只有一个函数就是vecadd()

      1. 首先产生llvm::Functionllvm::BasicBlock对象,分别对应函数与基本块。前面在loewr()函数中将函数的名为global_symbol的属性设为相应的函数名(如vecadd)。这里将该属性取出,作为生成函数的链接时的symbol。
      2. 通过VisitStmt()函数遍历IRModule中的各节点并转为LLVM中对应的数据结构,生成LLVM IR。这是最关键的一步了。前面费了老大劲构建起的TIR主要就是为了这里的转换。举例来说,对于ForNode就会调用CodeGenLLVM::VisitStmt_(ForNode *op)函数。它继而会调用CreateSerialFor()函数来产生相应的LLVM IR。在优化pass中的MakePackedAPImake_packed_api.cc)会添加一个AttrStmt,它对应一个值为目标函数名加_compute_后缀的compute_scope。这样,在code generation时 CodeGenCPU::CreateComputeScope()函数(为什么加compute_scope在该函数的注释中有提到)被调用。因此,最终的binary(可通过fadd.export_library("vecadd.so")语句导出)中大概会是这个样子:
        在这里插入图片描述
    6. AddMainFunction()函数设置主函数。如上面的例子中只有一个函数vecadd(),它也是主函数。这个symbol会放在runtime::symbol::tvm_module_main(即__tvm_main__)这个全局变量中。我们可以拿编译好binary验证这一点。用objdump命令dump导出的so文件,可以看到如下这段。如果将里边的0x766563616464的16进制转为ASCII,就是主函数的symbol名:vecadd。

    0000000000003c87 <__tvm_main__>:    
        3c87:   76 65                   jbe    3cee <__GNU_EH_FRAME_HDR+0x5e>
        3c89:   63 61 64                movslq 0x64(%rcx),%esp
        3c8c:   64                      fs     
    
    1. 最后,调用CodeGenCPU::Finish()函数将LLVM IR生成后端代码。它实际调用CodeGenLLVM::Finish()函数,它会调用CodeGenLLVM::Finish()函数。它主要调用CodeGenCPU::AddStartupFunction()函数和CodeGenLLVM::Optimize()函数。前者创建_tvm_module_startup函数,然后将一些需要启动时调用的函数填入。后者主要利用LLVM pass做一些优化。主要是向量化和函数内联。llvm中两种自动向量化。具体可参见Auto-Vectorization in LLVM

    其实,到这里编译还没有完全结束,只是构建好了LLVM的module。到这里,剩下的事情就是交给LLVM来编译生成可执行的binary了。真正生成可执行的binary是在第一次运行时通过LazyInitJIT()函数完成。 运行时会调用到LLVMModuleNode::GetFunction()函数。当它发现还未生成可执行binary时,会调用LazyInitJIT()函数。该函数通过llvm::ExecutionEngine将前面产生的llvm::Module编译成真正的(能在机器上跑的)binary。然后GetFunctionAddr()函数从中获得相应的函数指针,用于执行。

    展开全文
  • TVM46y5Y_Pi Network_v1.30.3三剑客社区.apk
  • tvm-financejs 常用财务功能库,用于货币时间价值计算。 介绍 该库是用于货币时间价值计算的一组通用财务公式。 公式被设计为具有与相同名称的Microsoft Excel财务公式相同的输入样式,并应提供基本相同的返回值。 ...
  • 初探TVM--TVM优化resnet50

    2021-10-30 10:33:25
    测试用TVM编译出的resnet50在CPU上的效果测试resnet50在CPU上的效果编译后的resnet50模型图像预处理运行编译后的模型查看输出结果resnet50自动调优模型调优 auto-tune编译调优过的模型 测试resnet50在CPU上的效果 ...

    测试resnet50在CPU上的效果

    如果直接点开了这篇,可能你会不知道编译过的模型是咋来的,戳这里。再回顾一下,编译过的模型会被压缩后存在一个tar压缩包里面。首先解压出来他:

    mkdir model
    tar -xvf resnet50-v2-7-tvm.tar -C model
    ls model
    

    你会看到model里面有三个文件:

    • mod.so 这个其实就是模型,只不过被编译为c++共享库,TVM的runtime会加载并调用它
    • mod.params 包含模型的预训练数据
    • mod.json 表示relay计算图的文本文件

    这些东西可以直接被你的应用加载,模型可以通过TVM的runtime API调用。

    编译后的resnet50模型

    我们已经编译出了模型模块,现在需要测试一下效果。测试使用tvm的runtime api,当然tvmc里面集成了它。使用时,我们需要准备:

    1. 编译过的模型,刚编出来,热乎的
    2. 一张输入的图片

    每个模型都会有期望的输入尺寸,数据类型,数据格式等等,因此对于一张图片,通常需要对齐进行预处理或者后处理。tvmc接受numpy的.npz文件,可以让我们简单的使用。我很喜欢猫子,这里就跟tvm教程里一样,就用这个猫子的照片了。
    请添加图片描述

    图像预处理

    对于resnet50,图像需要使用ImageNet的格式,下面放上一个pre-processing和post-processing的例子。在做前后处理的时候,需要使用pillow模块,如果没有的话,可以这样安装pip3 install pillow

    #!python ./preprocess.py
    from tvm.contrib.download import download_testdata
    from PIL import Image
    import numpy as np
    
    # if you have problem of download,just use images above
    img_url = "https://s3.amazonaws.com/model-server/inputs/kitten.jpg"
    img_path = download_testdata(img_url, "imagenet_cat.png", module="data")
    
    # Resize it to 224x224
    resized_image = Image.open(img_path).resize((224, 224))
    img_data = np.asarray(resized_image).astype("float32")
    
    # ONNX expects NCHW input, so convert the array
    img_data = np.transpose(img_data, (2, 0, 1))
    
    # Normalize according to ImageNet
    imagenet_mean = np.array([0.485, 0.456, 0.406])
    imagenet_stddev = np.array([0.229, 0.224, 0.225])
    norm_img_data = np.zeros(img_data.shape).astype("float32")
    for i in range(img_data.shape[0]):
            norm_img_data[i, :, :] = (img_data[i, :, :] / 255 - imagenet_mean[i]) / imagenet_stddev[i]
    
    # Add batch dimension
    img_data = np.expand_dims(norm_img_data, axis=0)
    
    # Save to .npz (outputs imagenet_cat.npz)
    np.savez("imagenet_cat", data=img_data)
    

    运行编译后的模型

    有了编译后的模型和转换后的图片,我们就可以测试模型的效果了:

    python -m tvm.driver.tvmc run \
    --inputs imagenet_cat.npz \
    --output predictions.npz \
    resnet50-v2-7-tvm.tar
    

    在tar文件包里面,有编译后的模型运行时库,tvmc封装了tvm的runtime接口,运行后,tvmc会给出一个预测结果的.npz文件。在这个例子中,运行模型的编译模型的机器为同一个平台,但你也可以使用RPC中提供的平台运算测试,通过python -m tvm.driver.tvmc run --help查看RPC使用的方式。

    查看输出结果

    其实每个模型都有自己的输出tensor格式,我们这里可以下载一个resnet50的输出查找表格,从中提取信息,并打印输出。这里会用到一个后处理的脚本:

    #!python ./postprocess.py
    import os.path
    import numpy as np
    
    from scipy.special import softmax
    
    from tvm.contrib.download import download_testdata
    
    # Download a list of labels
    labels_url = "https://s3.amazonaws.com/onnx-model-zoo/synset.txt"
    labels_path = download_testdata(labels_url, "synset.txt", module="data")
    
    with open(labels_path, "r") as f:
        labels = [l.rstrip() for l in f]
    
    output_file = "predictions.npz"
    
    # Open the output and read the output tensor
    if os.path.exists(output_file):
        with np.load(output_file) as data:
            scores = softmax(data["output_0"])
            scores = np.squeeze(scores)
            ranks = np.argsort(scores)[::-1]
    
            for rank in ranks[0:5]:
                print("class='%s' with probability=%f" % (labels[rank], scores[rank]))
    

    运行后可以拿到如下结果;

    class='n02123045 tabby, tabby cat' with probability=0.610552
    class='n02123159 tiger cat' with probability=0.367179
    class='n02124075 Egyptian cat' with probability=0.019365
    class='n02129604 tiger, Panthera tigris' with probability=0.001273
    class='n04040759 radiator' with probability=0.000261
    

    预测的top5全部都是不同种类的猫虎豹。

    resnet50自动调优

    上述模型仅仅完成基础的编译,并未加入任何与目标平台相关的调优工作,我们使用tvmc可以对模型根据目标平台特性,做自动调优。在一些情况下,我们其实不清楚在平台上使用哪些优化策略会比较好,auto-tune模块可以帮助我创建一个调优的搜索空间,并且进行性能调优。这里的tune并不是模型训练时的fine-tune,这里不改变模型的预测精度,仅仅是对目标平台的运行时速度做调优。tvm提供多个调优调度的模板,在目标平台中选出最优的那个,调优也可以通过tvmc实现。在最简单的调优模式中,tvmc需要我们给出:

    • 目标平台
    • 调优输出文件
    • 模型

    模型调优 auto-tune

    在下面的命令可以完成一次调优:

    python -m tvm.driver.tvmc tune \
    --target "llvm" \
    --output resnet50-v2-7-autotuner_records.json \
    resnet50-v2-7.onnx
    

    在这个例子中,我们可以写明我们的目标平台架构,例如cpu的skylake架构,你可以通过 – target llvm mcpu=skylake。这样auto-tune模块就可以找到更适用的算子优化组合。auto-tune模块优化在模型中的各个算子子图,每个子图都有一个优化调度的搜索空间,auto-tune会找出最佳的搜索结果1

    在这个模型中,得到如下tuning结果:

    [Task  1/25]  Current/Best:  201.89/ 220.20 GFLOPS | Progress: (40/40) | 20.89 s Done.
    [Task  2/25]  Current/Best:  163.58/ 188.97 GFLOPS | Progress: (40/40) | 10.87 s Done.
    [Task  3/25]  Current/Best:  107.46/ 232.10 GFLOPS | Progress: (40/40) | 13.03 s Done.
    [Task  4/25]  Current/Best:   86.44/ 216.15 GFLOPS | Progress: (40/40) | 18.86 s Done.
    [Task  5/25]  Current/Best:  158.27/ 225.16 GFLOPS | Progress: (40/40) | 12.15 s Done.
    [Task  6/25]  Current/Best:  127.20/ 204.70 GFLOPS | Progress: (40/40) | 15.10 s Done.
    [Task  7/25]  Current/Best:  181.57/ 211.43 GFLOPS | Progress: (40/40) | 12.19 s Done.
    [Task  8/25]  Current/Best:  136.64/ 205.29 GFLOPS | Progress: (40/40) | 21.14 s Done.
    [Task  9/25]  Current/Best:  235.01/ 235.01 GFLOPS | Progress: (40/40) | 20.43 s Done.
    [Task 10/25]  Current/Best:  121.53/ 209.96 GFLOPS | Progress: (40/40) | 12.77 s Done.
    [Task 11/25]  Current/Best:  214.43/ 214.43 GFLOPS | Progress: (40/40) | 12.11 s Done.
    [Task 12/25]  Current/Best:  147.68/ 206.92 GFLOPS | Progress: (40/40) | 14.60 s Done.
    [Task 13/25]  Current/Best:  112.67/ 214.38 GFLOPS | Progress: (40/40) | 13.11 s Done.
    [Task 14/25]  Current/Best:  168.30/ 216.40 GFLOPS | Progress: (40/40) | 20.96 s Done.
    [Task 15/25]  Current/Best:   68.19/ 221.29 GFLOPS | Progress: (40/40) | 20.52 s Done.
    [Task 16/25]  Current/Best:  140.36/ 194.06 GFLOPS | Progress: (40/40) | 13.67 s Done.
    [Task 17/25]  Current/Best:  104.27/ 219.83 GFLOPS | Progress: (40/40) | 13.28 s Done.
    [Task 18/25]  Current/Best:  102.57/ 214.91 GFLOPS | Progress: (40/40) | 17.39 s Done.
    [Task 19/25]  Current/Best:  141.44/ 223.53 GFLOPS | Progress: (40/40) | 14.55 s Done.
    [Task 20/25]  Current/Best:   62.98/ 202.96 GFLOPS | Progress: (40/40) | 21.22 s Done.
    [Task 21/25]  Current/Best:  100.66/ 218.22 GFLOPS | Progress: (40/40) | 19.84 s Done.
    [Task 22/25]  Current/Best:  162.85/ 215.54 GFLOPS | Progress: (40/40) | 12.81 s Done.
    [Task 23/25]  Current/Best:  122.70/ 231.81 GFLOPS | Progress: (40/40) | 16.73 s Done.
    [Task 24/25]  Current/Best:   23.29/  74.45 GFLOPS | Progress: (40/40) | 20.90 s Done.
    [Task 25/25]  Current/Best:    8.61/  85.78 GFLOPS | Progress: (40/40) | 20.67 s Done.
    

    调优的时间有时会比较长,所以tvmc提供其他选项供大家控制运算时间(--repeat或者--number)等。

    编译调优过的模型

    调优后的数据记录在resnet50-v2-7-autotuner_records.json文件中,这个文件可以在将来用作:

    • 编译优化后的模型 tvmc tune --tuning-records
    • 直接使用用作后续的进一步调优

    编译器会利用调优后的记录,在目标平台生成高效代码,可以用tvmc compile --tuning-records完成,也可以查看tvmc compile --help使用。收集过优化的数据,我们可以重新编译一遍模型了,:

    python -m tvm.driver.tvmc compile \
    --target "llvm" \
    --tuning-records resnet50-v2-7-autotuner_records.json  \
    --output resnet50-v2-7-tvm_autotuned.tar \
    resnet50-v2-7.onnx
    

    验证下优化的模型的预测结果:

    python -m tvm.driver.tvmc run \
    --inputs imagenet_cat.npz \
    --output predictions.npz \
    resnet50-v2-7-tvm_autotuned.tar
    
    python postprocess.py
    

    可以看到是相同的输出:

    class='n02123045 tabby, tabby cat' with probability=0.610552
    class='n02123159 tiger cat' with probability=0.367179
    class='n02124075 Egyptian cat' with probability=0.019365
    class='n02129604 tiger, Panthera tigris' with probability=0.001273
    class='n04040759 radiator' with probability=0.000261
    

    tvmc也提供了对比推理的运算时间的工具:

    python -m tvm.driver.tvmc run \
    --inputs imagenet_cat.npz \
    --output predictions.npz  \
    --print-time \
    --repeat 100 \
    resnet50-v2-7-tvm_autotuned.tar
    
    # Execution time summary:
    # mean (ms)   max (ms)   min (ms)   std (ms) 
    #   19.19      99.95      16.60       9.33 
    
    python -m tvm.driver.tvmc run \
    --inputs imagenet_cat.npz \
    --output predictions.npz  \
    --print-time \
    --repeat 100 \
    resnet50-v2-7-tvm.tar
    
    # Execution time summary:
    # mean (ms)   max (ms)   min (ms)   std (ms) 
    #   22.93      150.05     21.02      12.93
    

    可以看到,在我的这个服务器上面,时间提升了3~4ms,还算有点儿效果。撒花


    1. 搜索算法使用xgboost grid,但是也有别的算法可供选择,使用tvmc tune --help查看使用细节 ↩︎

    展开全文
  • TVM安装

    2020-02-16 12:49:09
    为什么选择TVM 为提升深度学习模型的推理效率,设备平台制造商针对自己的平台推出优化的推理引擎,例如NAVIDA的tensorRT,Intel的OpenVINO,Tencent针对移动端应用推出NCNN等。目前,深度学习模型应用广泛,在服务端...

    为什么选择TVM

    为提升深度学习模型的推理效率,设备平台制造商针对自己的平台推出优化的推理引擎,例如NAVIDA的tensorRT,Intel的OpenVINO,Tencent针对移动端应用推出NCNN等。目前,深度学习模型应用广泛,在服务端和移动端都有应用,甚至于特殊的嵌入式场景想,它们都有加速模型推理的需求。
      TVM介是从深度学习编译器的角度来做推理引擎,目前技术领域还比较新,具体技术细节以后有机会会深入学习,这里主要想体验一下使用TVM做深度模型推理,重点是推理效率的提升,所以尝试安装下TVM测试下。

    相关版本

    ```bash
    gcc版本 6.4.0
    cmake 3.16.4
    llvm 我选择的版本为llvm-9.0.0
    ```
    
    • gcc 安装

      • 查看系统版本
        cat /etc/redhat-release
        
      • 查看默认的gcc版本
        gcc --version
        
      • 查看默认动态库
        strings /usr/lib64/libstdc++.so.6 | grep GLIBC
        
        GLIBCXX_3.4
        GLIBCXX_3.4.1
        GLIBCXX_3.4.2
        GLIBCXX_3.4.3
        GLIBCXX_3.4.4
        GLIBCXX_3.4.5
        GLIBCXX_3.4.6
        GLIBCXX_3.4.7
        GLIBCXX_3.4.8
        GLIBCXX_3.4.9
        GLIBCXX_3.4.10
        GLIBCXX_3.4.11
        GLIBCXX_3.4.12
        GLIBCXX_3.4.13
        GLIBCXX_3.4.14
        GLIBCXX_3.4.15
        GLIBCXX_3.4.16
        GLIBCXX_3.4.17
        GLIBCXX_3.4.18
        GLIBCXX_3.4.19
        GLIBC_2.3
        GLIBC_2.2.5
        GLIBC_2.14
        GLIBC_2.4
        GLIBC_2.3.2
        GLIBCXX_DEBUG_MESSAGE_LENGTH
        
      • 下载gcc源码
        cd /usr/local/src/
        wget http://ftp.gnu.org/gnu/gcc/gcc-6.4.0/gcc-6.4.0.tar.xz
        tar -xf gcc-6.4.0.tar.xz -C /usr/src
        cd /usr/src/gcc-6.4.0
        # 自动下载mpfr-2.4.2.tar.bz2、gmp-4.3.2.tar.bz2、mpc-0.8.1.tar.gz和isl-0.15.tar.bz2软件包
        ./contrib/download_prerequisites
        ./configure -enable-checking=release -enable-languages=c,c++ -disable-multilib
        make -j4
        make install
        
      • 配置gcc
        find / -name "libstdc++.so*"
        # /usr/src/gcc-6.4.0/stage1-x86_64-pc-linux-gnu/libstdc++-v3/src/.libs/libstdc++.so.6.0.22
        
        cd /usr/lib64
        cp /usr/src/gcc-6.4.0/stage1-x86_64-pc-linux-gnu/libstdc++-v3/src/.libs/libstdc++.so.6.0.22 libstdc++.so.6.0.22
        mv libstdc++.so.6 libstdc++.so.6.old
        ln -sv libstdc++.so.6.0.22 libstdc++.so.6
        #######
        strings /usr/lib64/libstdc++.so.6 | grep GLIBC
        
        GLIBCXX_3.4
        GLIBCXX_3.4.1
        GLIBCXX_3.4.2
        GLIBCXX_3.4.3
        GLIBCXX_3.4.4
        GLIBCXX_3.4.5
        GLIBCXX_3.4.6
        GLIBCXX_3.4.7
        GLIBCXX_3.4.8
        GLIBCXX_3.4.9
        GLIBCXX_3.4.10
        GLIBCXX_3.4.11
        GLIBCXX_3.4.12
        GLIBCXX_3.4.13
        GLIBCXX_3.4.14
        GLIBCXX_3.4.15
        GLIBCXX_3.4.16
        GLIBCXX_3.4.17
        GLIBCXX_3.4.18
        GLIBCXX_3.4.19
        GLIBCXX_3.4.20
        GLIBCXX_3.4.21
        GLIBCXX_3.4.22
        GLIBC_2.3
        GLIBC_2.2.5
        GLIBC_2.14
        GLIBC_2.17
        GLIBC_2.3.2
        GLIBCXX_DEBUG_MESSAGE_LENGTH
        
    • cmake 安装
      https://cmake.org/download/

      # 解压Cmake压缩包
      tar -zxvf cmake-3.16.4.tar.gz
      # 进入Cmake文件夹
      cd cmake-3.16.4
      # 运行bootstrap
      ./bootstrap
      # 编译gmake
      gmake
      # 进行安装Cmake
      gmake install
      

    Step1 llvm安装

    llvm源码下载
    llvm安装教程
    在这里插入图片描述
    下载 Sources下除LLVM Test Suite外的所有源码

    • 一键下载
      # LLVM 源码下载
      wget https://releases.llvm.org/9.0.0/clang-tools-extra-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/compiler-rt-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/libunwind-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/libcxxabi-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/libcxx-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/openmp-9.0.0.src.tar.xz 
      wget https://releases.llvm.org/9.0.0/lldb-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/polly-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/lld-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/cfe-9.0.0.src.tar.xz
      wget https://releases.llvm.org/9.0.0/llvm-9.0.0.src.tar.xz
      # 放在llvm/tools 目录下的组件
      tar Jxvf llvm-9.0.0.src.tar.xz
      mv llvm-9.0.0.src llvm
      
      tar Jxvf cfe-9.0.0.src.tar.xz
      mv cfe-9.0.0.src clang
      mv clang llvm/tools
      
      tar Jxvf lld-9.0.0.src.tar.xz
      mv lld-9.0.0.src lld
      mv lld llvm/tools
      
      tar Jxvf polly-9.0.0.src.tar.xz
      mv polly-9.0.0.src polly
      mv polly llvm/tools
      
      tar Jxvf lldb-9.0.0.src.tar.xz
      mv lldb-9.0.0.src lldb
      mv lldb llvm/tools
      
      # 放在 llvm/projects 目录下的组件
      tar Jxvf openmp-9.0.0.src.tar.xz 
      mv openmp-9.0.0.src openmp
      mv openmp llvm/projects
      
      tar Jxvf libcxx-9.0.0.src.tar.xz
      mv libcxx-9.0.0.src libcxx
      mv libcxx llvm/projects
      
      tar Jxvf libcxxabi-9.0.0.src.tar.xz
      mv libcxxabi-9.0.0.src libcxxabi
      mv libcxxabi llvm/projects
      
      tar Jxvf libunwind-9.0.0.src.tar.xz
      mv libunwind-9.0.0.src libunwind
      mv libunwind llvm/projects
      
      
      tar Jxvf compiler-rt-9.0.0.src.tar.xz
      mv compiler-rt-9.0.0.src compiler-rt
      mv compiler-rt llvm/projects
      # clang 的子组件
      tar Jxvf clang-tools-extra-9.0.0.src.tar.xz
      mv clang-tools-extra-9.0.0.src extra
      mv extra llvm/tools/clang/tools
      # 删除下载的tar.xz
      rm -rf *.tar.xz
      
      llvm 为了防止污染源码,不能在源码目录直接安装,在 llvm 同级目录下新建一个 build 目录,进入再安装
      cmake -G "Unix Makefiles"  \
      -DLLVM_ENABLE_ASSERTIONS=On  \
      -DCMAKE_BUILD_TYPE=Release \
      -DCMAKE_C_COMPILER=/usr/local/bin/gcc \
      -DCMAKE_CXX_COMPILER=/usr/local/bin/c++ \
      -DLIBCXXABI_LIBC_INCLUDES= /path/llvm/projects/libcxxabi/include  \
      -DLIBCXX_LIBCXX_INCLUDES= /path/llvm/projects/libcxx/include \
      ../llvm
      
      make -j4
      make install
      
      如果出现以下报错:
      CMake Error at tools/lldb/scripts/cmake_install.cmake:41 (file):
      file INSTALL cannot find "/usr/local/src/llvm/llvm/build/./lib/python2.7":
      No such file or directory.
      Call Stack (most recent call first):
      tools/lldb/cmake_install.cmake:50 (include)
      tools/cmake_install.cmake:51 (include)
      cmake_install.cmake:68 (include)
      
      执行
      cp -r lib64/python2.7 lib/python2.7
      
      如果出现以下报错:
      CMake Error at /usr/local/share/cmake-3.16/Modules/FindPackageHandleStandardArgs.cmake:146 (message):
      Could NOT find SWIG (missing: SWIG_EXECUTABLE SWIG_DIR)
      Call Stack (most recent call first):
      /usr/local/share/cmake-3.16/Modules/FindPackageHandleStandardArgs.cmake:393 (_FPHSA_FAILURE_MESSAGE)
      /usr/local/share/cmake-3.16/Modules/FindSWIG.cmake:64 (FIND_PACKAGE_HANDLE_STANDARD_ARGS)
      tools/lldb/scripts/CMakeLists.txt:16 (find_package)
      
      执行
      yum install swig
      
      如果报以下错
      CMake Error at /usr/local/share/cmake-3.16/Modules/FindPackageHandleStandardArgs.cmake:146 (message):
      Could NOT find LibEdit (missing: libedit_INCLUDE_DIRS libedit_LIBRARIES)
      Call Stack (most recent call first):
      /usr/local/share/cmake-3.16/Modules/FindPackageHandleStandardArgs.cmake:393 (_FPHSA_FAILURE_MESSAGE)
      tools/lldb/cmake/modules/FindLibEdit.cmake:54 (find_package_handle_standard_args)
      tools/lldb/cmake/modules/LLDBConfig.cmake:94 (find_package)
      tools/lldb/CMakeLists.txt:20 (include)
      
      执行
      yum install libedit.i686 libedit-devel.i686
      

    Step2 TVM安装

    tvm Install from Source

    git clone --recursive https://github.com/apache/incubator-tvm tvm
    git submodule init
    git submodule update
    
    • 定制config.cmake
    cd tvm $$ mkdir build
    cp cmake/config.cmake build
    cd build
    编辑build/config.cmake文件,里面有一些功能开关,这些配置有:
    
    USE_CUDA,NVIDIA的GPU计算;
    USE_ROCM,通用的GPU计算,AMD提出,目的很显然...;
    USE_SDACCEL,FPGA计算;
    USE_AOCL,Intel FPGA SDK for OpenCL (AOCL) runtime;
    USE_OPENCL,异构平台编写程序的框架,异构平台可由CPU、GPU、DSP、FPGA或其他类型的处理器与硬件加速器所组成;
    USE_METAL,iOS上的GPU计算;
    USE_VULKAN,新一代的openGL,Android 7.x开始支持(iOS不支持,因为有自己的metal2);
    USE_OPENGL,2D/3D渲染库标准,显卡厂家负责实现和支持;
    USE_SGX, Intel SGX ; 
    USE_RPC,远程调用,电脑和手机可以通过网络联调;
    USE_STACKVM_RUNTIME,embed stackvm into the runtime;
    USE_GRAPH_RUNTIME,enable tiny embedded graph runtime;
    USE_GRAPH_RUNTIME_DEBUG,enable additional graph debug functions;
    USE_LLVM,llvm support;
    USE_BLAS,API标准,规范发布基础线性代数操作的数值库(如矢量或矩阵乘法),不同的实现有openblas, mkl, atlas, apple
    USE_RANDOM,contrib.random运行时;
    USE_NNPACK,
    USE_CUDNN,
    USE_CUBLAS,
    USE_MIOPEN,
    USE_MPS,
    USE_ROCBLAS,
    USE_SORT,使用contrib sort;
    USE_ANTLR,
    USE_VTA_TSIM,
    USE_RELAY_DEBUG,Relay debug模式
    
    cmake .. -DCMAKE_INSTALL_PREFIX=./install  //如果需要gdb跟踪源码的话需要加-DCMAKE_BUILD_TYPE=Debug
    make -j4
    make install
    

    error:

    git submodule init
    perl: warning: Setting locale failed.
    perl: warning: Please check that your locale settings:
            LANGUAGE = (unset),
            LC_ALL = (unset),
            LANG = "zh_CN.utf8"
        are supported and installed on your system.
    perl: warning: Falling back to the standard locale ("C").
    

    解决方法:

    export LANGUAGE=en_US.UTF-8
    export LC_ALL=en_US.UTF-8
    export LANG=en_US.UTF-8
    

    tvm系统配置

    vim ~/.bashrc
    
    export TVM_PATH=/yourpath/to/tvm  # 你自己tvm
    export PYTHONPATH=$TVM_PATH/python:$TVM_PATH/topi/python:$TVM_PATH/nnvm/python:${PYTHONPATH}
    
    source ~/.bashrc
    

    安装TVM python bindings by setup.py

    cd tvm
    cd python
    python3 setup.py install --user
    cd ..
    cd topi/python
    python setup.py install --user
    

    测试

    python3
    Python 3.6.5 |Anaconda, Inc.| (default, Apr 29 2018, 16:14:56) 
    [GCC 7.2.0] on linux
    Type "help", "copyright", "credits" or "license" for more information.
    >>> import tvm
    >>> print(tvm.__version__)
    0.7.dev0
    >>> 
    

    Reference

    PyTorch模型转TVM模型全流程记录
    安装TVM
    TVM学习(1)–搭建环境
    --------------------------------------------------------- END --------------------------------------------------------

    展开全文
  • TVM

    2020-12-14 13:40:51
    TVM简介: https://blog.csdn.net/tlzhatao/article/details/93630910 TVM: An Automated End-to-End Optimizing Compiler for Deep Learning

    TVM简介: https://blog.csdn.net/tlzhatao/article/details/93630910

    展开全文
  • tvm:图灵虚拟机 没人关心的图灵机模拟器。 正在进行中-我可能会也可能不会完成。
  • TVM在行动 在我探究深度学习框架的令人难以置信的爆炸以及如何将它们组合在一起的过程中,此存储库包含我的笔记,TVM堆栈的教程材料(源代码)。 抽象的 TensorFlow,MXNet,Caffe和PyTorch等可扩展框架针对各种服务...
  • 0x03.2 本地编译以Ubuntu为例 如果有修改TVM源码或者给TVM贡献的需求,可以本地编译TVM,以Ubuntu为例编译和配置的流程如下: git clone --recursive https://github.com/apache/tvm tvm cd tvm mkdir build ...
  • TVM论文阅读

    2021-01-29 22:06:02
    TVM论文阅读1 介绍1.1 优化的基本挑战1.2 TVM:一个端到端优化堆栈2 优化计算图2.1 计算图2.2 操作符融合2.3 数据布局转换2.4 计算图级别优化的限制3 优化张量操作3.1 Tensor表达式语言3.2 调度空间3.3 协作式嵌套...
  • TVM 架构设计

    千次阅读 2020-12-08 09:10:33
    TVM 架构设计 本文面向希望了解TVM体系结构和/或积极参与项目开发的开发人员。 主要内容如下: 示例编译流程概述了TVM将模型的高级概念转换为可部署模块的步骤。 逻辑架构组件部分描述逻辑组件。针对每个逻辑组件,...
  • 初识 TVM

    千次阅读 2019-11-15 22:15:24
    笔者也是最近偶然的机会才开始接触TVM,使用过后发现,经过auto-tuning后的TVM模型在速度是竟然超过了TensorRT,并且笔者使用的是MXNet框架,TVM对MXNet绝对的友好,对于Pytorch等模型,可以使用ONNX,操作也一样简单...
  • TVM Pass 总结

    千次阅读 2020-07-14 11:50:54
    介绍 tvm.relay.transform
  • 部署TVM Runtime

    2021-05-06 05:30:43
    本文主要介绍如何在开发板上部署TVM Runtime, 在本地机器安装完整的TVM(包含了TVM Runtime以及编译功能), 并且使用一个简单的远程调用例子测试是否部署成功。本地机器使用的是Linux操作系统,开发板使用的是预装的...

空空如也

空空如也

1 2 3 4 5 ... 20
收藏数 4,109
精华内容 1,643
关键字:

TVM