Tensorflow XLA Service Buffer优化详解

下图是Tensorflow 架构图,以及XLA在Tensorflow中的位置

通过层层转换(参考Tensorflow XLA Service 详解 I),Graph在进入XLA Service前已经被表达为HloModule的形式,而作为图编译器的核心,XLA Service就负责将HloModule表达的计算图编译为可以直接在不同硬件平台(Backend)执行的程序,而编译的核心,就是优化代码,包括设备无关的优化和设备相关的优化:

  1. 优化HloModule所表示的计算图,并将其转化为LLVM HLO
  2. 基于LLVM,生成硬件相关的二进制

作为通用的编译器框架,LLVM 会对LLVM HLO做大量的优化,生成高效的Target binary, 所以作为XLA的开发者,主要关注阶段1的优化: 设备无关的图优化算法。Compiler是XLA适配硬件的接口类,每个适配XLA的硬件都必须实现其中的方法。尤其是RunBackend(),参考Tensorflow XLA Service 详解 I 一文,该接口是进行图优化和编译的入口,统领整个优化和编译过程。

同时,XLA Service还实现了一组通用的优化方法(各种Schedule策略,各种Memory优化算法)供各个硬件平台的编译器使用,当然,主要是供给RunBackend()调用。以XLA的GPU平台的编译优化流程为例:

NVPTXCompiler::RunBackend()
  hlo_schedule = GpuHloSchedule::Build(*module, *stream_assignment, pointer_size_)
  BufferAssigner::Run(hlo_schedule->ConsumeHloOrdering()...)
  entry_computation->Accept(&ir_emitter)
  CompileToPtx() 
  CompilePtxOrGetCachedResult()

-1- 从XLA Service通用层中选择适合GPU的Schedule策略
-3- 基于Schedule策略,进行设备无关的Buffer优化,主要关注尽可能的减少Buffer的大小。注意,这里是设备无关的优化,是无法利用硬件Memory特性的。
-4- 将HloModule转化为LLVM IR
-5,6- 利用LLVM框架,将LLVM IR编译为二进制代码。

本文主要关注-3-,是XLA优化的核心。

对BufferAssigner::Run()进一步分解。

NVPTXCompiler::RunBackend()
  hlo_schedule = GpuHloSchedule::Build(*module, *stream_assignment, pointer_size_)
  //this analysis figures out which temp buffers are required to run the computation
  BufferAssigner::Run(hlo_schedule->ConsumeHloOrdering()...)
    assigner.CreateAssignment(HloModule, hlo_ordering, buffer_size)
      liveness = BufferLiveness::Run()
      assignment = new BufferAssignment(module, liveness, ...)
      set<LogicalBuffer*> colocated_buffers
      set<BufferAllocation::Index> colocated_allocations
      vector<ColocatedBufferSet> colocated_buffer_sets
      BuildColocatedBufferSets(&colocated_buffer_sets)
      colorer_(assignment->liveness())
      AssignColocatedBufferSets(colocated_buffer_sets, assignment, &colocated_buffers, &colocated_allocations);
      GatherComputationsByAllocationType(module, &thread_local_computations, &global_computations)
      for computation : global_computations:
        AssignBufferForComputation(computation, false, buffers_to_assign_sequentially)
      AssignBuffersWithSequentialOrdering(buffers_to_assign_sequentially, ,assignment)
      for computation : thread_local_computations:
        AssignBuffersForComputation()
      for buffer : assignment->liveness().maybe_live_out_buffers():
        if assignment->HasAllocation(buffer):
          assignment->GetMutableAssignedAllocation(buffer).set_mayby_live_out(true)
      assignment->CombineTempAllocations()
      return std::move(assignment)
  entry_computation->Accept(&ir_emitter)
  CompileToPtx() 
  CompilePtxOrGetCachedResult()

-6- 进行BufferLiveness分析,分析整个HloModule的LogicalBuffer的干涉关系,为后续优化提供依据
-11- BuildColocatedBufferSets, 依据Bufferliveness的分析,将所有的LogicalBuffer分为几个Bufferset,并进行初步的Set融合,每个Bufferset内
参照注释, colocated buffer sets, 每个set都是一组可以共享BufferAllocation的LogicalBuffer, 共享Allocation,意味着共享同一块物理内存(GPU的显存)
-12- colorer_ 缺省被赋值为BufferLiveness::DefaultColorer(), 所有的LogicalBuffer实例的color都会被设置为0
-13- AssignColocatedBufferSets, 为Bufferset分配BufferAllocation, 每一个LogicalBufferSet 与其关联, 这里用到了buffer_size_, 这个函数是判断一个LogicalBuffer大小, LogicalBuffer的大小要和相应的Allocation一样, 具体可以参考tf2xla/while_op.cc tf2xla/if_op.cc xla/client/builder.cc kConditional代码,可以看到明显的要求各个body的Shape要一致。通过TEST用例也能确认
-14- GatherComputationsByAllocationType,根据内含的LogicalBuffer的属性,将Allocation分为global和thread local两类,这部分是理解显存优化的关键,后文详细
-16- AssignBufferForComputation,关联Allocation和XlaComputation,此调用点只针对global,temp buffer被收集到buffers_to_assign_sequentially, 延后处理,

BufferLiveness::Run()

Continue reading

BERT Training By Example

BERT 模型的训练主要分为三个阶段:
1. Masked Language Model,基于对自然语言的语料打标签生成训练样本,将无监督学习转换为监督学习
2. Pretraining,模型训练的核心过程,生成能理解语言的模型
3. Fine-tuning,基于对模型的理解,完成特定的任务。

Masked Language Model

Masked Language Model 的任务是将无标签的语料转换为有标签的语料,进而将无监督学习转化为监督学习。这个过程需要准备两份材料:原始语料词汇表
原始语料就是一系列符合语法的句子或段落。
词汇表维持token-id的映射,除了一门语言自身的词汇,BERT还要求词汇表包含[PAD], [UNK], [CLS], [SEP]和[MASK]。

Pretraining

Continue reading

为Tensorflow 添加新Op

Tensorflow内部添加op有3种方式:
1. 不修改Tensorflow源码, 在运行环境里通过Tensorflow包中的.h和.so直接扩展, 无需Tensorflow编译环境, 二次开发首选
2. 修改Tensorflow源码, 利用源码中的User-supplied Op接口直接将新Op添加到Tensorflow安装包中, 关注Op实现, 无需调试Tensorflow的Bazel编译框架, 但个人认为比较鸡肋
3. 修改Tensorflow源码, 以Standard Op的方式将Op添加在Tensorflow包中, 将新Op提交到社区需要这种方式. 需要实现Op, 同时需要修改Tensorflow的Bazel编译框架, 由于Tensorflow使用了自动生成代码的技术, 后者调试比较困难,谨慎使用

本文主要介绍这3种添加TfOp的方式, 最后, 简要介绍添加XlaOp的方法. 对于需要修改源码的方法2和方法3, 在Tensorflow内部均需要修改python接口和tfop两层, 而每一种又有两种方式:User-supplied Op 和Standard Op, 所以一共是2*2 = 4 种, 简化起见, 本文分别用用user-supplied op实现了XZeroOutOp(输入清零后输出), 用Standard Op 实现了XDoNothingOp(输入透传给输出), 在实际操作中, tensorflow内部这两层使用自动生成代码技术进行了解耦, python层和tf层使用不同的方式完全没问题.

Customized Op

Horovod就是使用Cutomized Op这种方式的典型框架, 基于Tensorflow 接口, MPI和NCCL接口, Horovod在不修改Tensorflow源码的前提下实现了Ring-allreduce 的Op. 使用这种方式Op实现的方式与Tensorflow内部的Standard Op并无区别, 关键是编译的时候要include相关的头文件:”/usr/lib/python2.7/site-packages/tensorflow/include/tensorflow/core/framework/”, 以及链接相应的库:”/usr/lib/python2.7/site-packages/tensorflow/libtensorflow_framework.so”

此外需要注意, Tensorflow框架会自动将驼峰命名转化为下划线命名, 即名为”HorovodAllreduce”的Op, 在生成的so中, 对应”horovod_allreduce”, 感兴趣的同学可以参考Tensorflow自动生成py的代码.

User-supplied Op

Continue reading

Tensorflow-Horovod安装部署checklist

基础环境 GPU 检查多机中每个节点的GPU是否归属同一型号, horovod后台线程周期性处理提交的梯度, 如果由于GPU计算能力不同或网络延迟导致某次allreduce中某个提交者落后超过一个周期, 会发生使用错误的Tensor进行allreduce的情况, 通常, 这都会导致Tensor Shape Mismatch的错误, 笔者就层遇到过不小心同时使用M40和P40的情况, 在tensorflow1.8还能的运行, 但1.14就无法多机reduce, 定位了好久才发现是硬件不同
cuda 是否与驱动match:cuda兼容性检查下载, 随同一起安装的还有cupti等组件
cudnn 版本是否与cuda版本match: cudnn兼容性检查及下载
nccl 版本是否与cuda版本match: nccl兼容性检查及下载
mpi 其中libmpi.so要对外可见, 多机训练需要tensorflow编译的时候就要将MPI链接到lib中, 同时, MPI对于horovod也是必要依赖
Tensorflow tensorflow 选择的tensorflow 版本是否兼容系统cuda, 官方兼容列表
tensorboard tensorboard是否兼容tensorflow, tensorboard的版本号要与tensoflow保持一致
tensorflow-estimator tensorflow-estimator是否兼容了tensorflow, tensorflow 1.10.6 之后的特性, 是对底层API的封装, 需要与tensorflow版本号一致
bazel tensorflow对使用的bazel版本比较敏感, 如果没有修改bazelrc却出现bazelrc文件解析错误, 就要考虑下是不是bazel版本的问题, 参考tensorflow测试通过的bazel版本bazel下载
编译选项 tensorflow编译选项, 是否包含了cuda和nccl, 这两个缺省都是不安装的, 此处使用GPU, cuda为必须, 加之多机多卡, NCCL可以减少通信开销, 应该编译进tensorflow. 在1.13之前, nccl是随同config的过程一同配置, 1.14开始使用环境变量 + .bazelrc控制NCCL的安装, 具体的, 环境变量中需要export以下环境变量
export TF_NCCL_VERSION=2.4.2
export NCCL_HDR_PATH=/usr/local/nccl_2.4.2-1+cuda9.2_x86_64/include/
            
同时, 修改tensorflow/.bazelrc文件
build:nccl --define=with_nccl_support=true
#build:nccl --define=no_nccl_support=true
            
Horovod 环境变量 环境变量是否指定了libtensorflow_framework.so的路径以及NCCL路径信息
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/lib/python2.7/site-packages/tensorflow/
export HOROVOD_WITH_TENSORFLOW=1
export HOROVOD_WITHOUT_PYTORCH=1
export HOROVOD_GPU_ALLREDUCE=NCCL
export HOROVOD_NCCL_HOME=/usr/local/nccl_2.1.15-1+cuda9.0_x86_64
export HOROVOD_NCCL_INCLUDE=/usr/local/nccl_2.1.15-1+cuda9.0_x86_64/include/
export HOROVOD_NCCL_LIB=/usr/local/nccl_2.1.15-1+cuda9.0_x86_64/lib/

Tensorflow XLA Service 详解 II

本文主要介绍在XLA service阶段针对HloInstruction做的一些显存优化, 对于训练框架来说, 显存优化的工作至关重要, 主要是由于现阶段GPU+CUDA远没有CPU+Linux组合强大, 后者有完善的建立在虚拟内存基础上的内存管理机制, 内存的高效使用由linux kernel来负责, 即便物理内存不足, 还可以使用swap, 内存压缩等技术确保内存的高效供应, 而在GPU+CUDA里, 这方面的工作很大程度让渡给了程序员自己来搞定, GPU程序接触到的就是物理显存, 如果程序的显存申请超过显存容量, 整个程序就会直接coredump, 此外, 显存本身就集成在GPU板卡上, 无法像内存一样扩展, 而GPU本身造价昂贵, 最后, 在深度学习训练中, 大力出奇迹的现状下, 显存的消耗明显超过的摩尔定律, 这也加剧了显存供求关系的矛盾, 正式由于训练框架做了大量的优化, 才能让模型跑起来.

XLA Service的显存优化设计思想与tensorflow整体一样遵循”静态图”的设计: 先整体优化, 再落地实施. 其中, xla/service/buffer_assignment.cc 是整个显存优化的核心, 在1.14版本中, xla/service/支持两种后端: cpu和gpu, 纷纷针对两种backend有进一步的优化算法, 本文主要针对GPU的优化逻辑进行分析

核心文件

内存优化公共:
xla/service/buffer_assignment 内存优化核心文件
xla/service/buffer_liveness.cc  内存片生命周期分析
GPU相关:
xla/service/gpu/buffer_allocations.cc BufferAllocation的组合
xla/service/gpu/gpu_hlo_scheudle.cc   Hlo的处理顺序, 与显存的优化策略息息相关, 简单地说, 按照BFS并行执行的HloInstruction消耗的显存肯定大于所有的HloInstruction都顺序执行. 

核心抽象

HloSchedule

XxxSchedule是TF的代码风格, 类似的有更底层用于Thunk调度的ThunkSchedule, 以及Service提供的HloSchedule. HloSchedule中最重要的就是封装了HloOrdering.

HloSchedule

XLAService内存优化的本质是处理LoigicalBuffer和BufferAllocation之间的关系, 原则是使用尽可能少的BufferAllocation去承载尽可能多的LogicalBuffer, 而如何使用的更少, 就涉及到了对Hlo图的分析, 就涉及到了Ordering的问题, 使用不同策略生成Ordering, 直接影响两个LogicalBuffer之间的约束关系, 最简单的, 在图遍历中, 使用DFS和BFS的2种方式遍历会导致图上节点的内存依赖关系大有不同.

HloOrdering是描述HloInstruction加载序列的基类, 派生类有PredecessorHloOrdering, DependencyHloOrdering 和 SequentialHloOrdering, 其中, DependencyHloOrdering基于依赖关系, 所以可以并行, 性能更高, 但耗更多的内存, 而SequentialHloOrdering完全串行, 性能相对低, 但可以节约更多内存, 而 PredecessorHloOrdering 是个虚类, 需要子类进一步填充predecessors_, 这也是GPU后端使用的方式.不同的Ordering会影响内存的依赖关系, 进一步影响Launch到GPU后Kernel的执行序列. 

Continue reading

Tensorflow XLA Service 详解 I

compiler/aot/ 以AOT的方式将tf2xla/接入TF引擎
compiler/jit/以JIT的方式将tf2xla/接入TF引擎, 核心是9个优化器和3个tfop,其中XlaCompileOp调用tf2xla的“编译”入口完成功能封装,XlaRunOp调用xla/client完成“运行”功能。
compiler/tf2xla/对上提供xla_compiler.cc:XlaCompiler::CompileFunction()供jit:compile_fn()使用将cluster转化为XlaComputation。核心是利用xla/client提供的接口,实现对XlaOpKernel的“Symbolic Execution”功能。每个XlaOpKernel子类均做的以下工作: **从XlaOpKernelContext中取出XlaExpression或XlaOp, 调用xla/client/xla_buidler.h提供的方法完成计算, 将计算结果的XlaOp存入XlaKernelContext.**
compiler/xla/client/ 对上提供xla_builder.cc:Builder等供CompileFunction()使用,将Graph由Op表达转化为HloModuleProto:HloComputationProto:HloInstructionProto表达并保存在XlaComputation中。
对上提供local_client.cc:LocalClient::Compile(),作为编译入口供jit:BuildExecutable()使用,将已经得到的XlaComputation交给service并进一步编译为二进制。
对上提供local_client.cc:LocalExecutable::Run(),作为运行入口供jit/kernels/xla_ops.cc:XlaRunOp使用,通过Key找到相应的二进制交给service层处理
compiler/xla/service/ 对上提供local_service.cc:LocalService::BuildExecutable()供LocalClient::Compile()使用实现真正的编译,承接XlaComputation封装的HloProto, 将其转化为HloModule:HloComputation:HloInstruction表达, 对其进行优化之后, 使用LLVM后端将其编译为相应Executable后端的二进制代码
对上提供executable.cc:Executable::ExecuteOnStream()供LocalExecutable::Run()使用实现真正的执行二进制。

编译cubin

调用栈:

Continue reading

Tensorflow XLA Client | HloModuleProto 详解

compiler/aot/ 以AOT的方式将tf2xla/接入TF引擎
compiler/jit/以JIT的方式将tf2xla/接入TF引擎, 核心是9个优化器和3个tfop,其中XlaCompileOp调用tf2xla的“编译”入口完成功能封装,XlaRunOp调用xla/client完成“运行”功能。
compiler/tf2xla/对上提供xla_compiler.cc:XlaCompiler::CompileFunction()供jit:compile_fn()使用将cluster转化为XlaComputation。核心是利用xla/client提供的接口,实现对XlaOpKernel的“Symbolic Execution”功能。每个XlaOpKernel子类均做的以下工作: 从XlaOpKernelContext中取出XlaExpression或XlaOp, 调用xla/client/xla_buidler.h提供的方法完成计算, 将计算结果的XlaOp存入XlaKernelContext.**
compiler/xla/client/ 对上提供xla_builder.cc:Builder等供CompileFunction()使用,将Graph由Op表达转化为HloModuleProto:HloComputationProto:HloInstructionProto表达并保存在XlaComputation中。
对上提供local_client.cc:LocalClient::Compile(),作为编译入口供jit:BuildExecutable()使用,将已经得到的XlaComputation交给service并进一步编译为二进制。
对上提供local_client.cc:LocalExecutable::Run(),作为运行入口供jit/kernels/xla_ops.cc:XlaRunOp使用,通过Key找到相应的二进制交给service层处理
compiler/xla/service/ 对上提供local_service.cc:LocalService::BuildExecutable()供LocalClient::Compile()使用实现真正的编译,承接XlaComputation封装的HloProto, 将其转化为HloModule:HloComputation:HloInstruction表达, 对其进行优化之后, 使用LLVM后端将其编译为相应Executable后端的二进制代码
对上提供executable.cc:Executable::ExecuteOnStream()供LocalExecutable::Run()使用实现真正的执行二进制。

compiler/xla/client 向上为tf2xla/下的XlaOpKernel的实现提供支撑, 将上层请求转换为HloModule交给下层xla/service优化并编译.

接口上, client做上表中的三件事 , 实际上, 只有XlaOpKernel->HloProto在Client完成, 而对于另外两个编译和执行,都是在service中完成的,或者说,Client本质上是Service的Facade + Proxy,对繁琐的构造HloModuleProto的过程进行了封装,降低了XLA Service与上层模块的耦合

 client.h:Client    Client基类, 用于多态实现
 client_library.h:ClientLibarary  使用单例构造client_library对象,  用于检索/构造所需的Client实例
 lib/   同builder一起实现”Symbolic Execution”
 local_client.h:LocalClient, LocalExecutable JIT 使用的LocalClient定义, 是Service相关方法的Proxy
 xla_builder.h:XlaBuilder    提供接口用tf2xla使用实现”Symbolic Execution”, 是其中XlaBuilder::Build()是构造client构造HloModuleProto的核心方法。对于需要多个步骤完成初始化的类, 我们会使用Builder模式, 这就是个例子
 xla_computation.h:XlaComputationXlaComputation对象是对HloModuleProto的封装, 用于进一步二进制编译

职责1: 构造HloModuleProto

编译二进制之前首先要完成Graph表达方式的映射: Client之前的tf2xla的Graph由使用Op表达, Client之后的Service的Graph使用HloInstruction表达, Client负责完成这种转化, 具体地, 就是将Op转化为HloProto格式, 再交给Service解析为Hlo格式, 其中的HloProto就是封装在XlaComputation中. 所以, 这个过程可以看做是”编译”的准备工作. 在这个过程中, Graph, Cluster, XlaComputation, HloModuleProto, HloModule是一一对应的,都是”Program”的概念, 这一点从源码的ProgramShape等变量名中都可以体现

按照OOP的方式理解OOP的代码, 在说明Client是如果构造HloModuleProto之前, 来介绍几个概念, HloModuleProto-HloComputationProto- HloInstructionProto, 这3个概念是XLA Service提供给上层调用者(Here, XLA Client)的用于构造HloModule-HloComputation-HloInstruction的protobuf形式的”原材料”, Module-Computation-Instruction, 在概念上可以对应程序-函数-语句

  1. HloProto中的函数调用都是以”半inline“方式实现, inline的部分是指类似”a = add(x, y);”这样的语句, 会将add函数体整个copy到该语句的上下文, 在执行的时候调用之, 在HloProto中, 被调用的XlaComputation也会被copy到调用语句的上下文, “上下文”就是调用语句HloInstructionProto所处的XlaComputation(对应一个HloModuleProto)的XlaBuilder实例. “半”的部分是指被copy的XlaComputation并不是直接在调用语句处展开, 而是在调用语句的HloModuleProto中存储被调用的XlaComputation的Id, 执行的时候直接跳转执行.
  2. 类似多文件编程, 构造的过程可以有多个HloModuleProto, 但由于上一条所述, 最终所有的被调用函数均会被copy到根HloModuleProto的上下文, 最终只有也只需要根HloModuleProto就足以以描述整个程序的逻辑
  3. 一个C/C++程序有main作为入口, 对应到HloModuleProto就是entry_computation, 一个函数的第一条语句, 对应到HloComputationProto就是root_id

下面我们简要分析下代码, 既然Hlo是描述一个程序, 那么就需要解决一个基本的问题: 描述出函数->语句->函数->语句…这样的递归结构, OOP的Composite模式, 或者内核的kset-kobj结构, 都是解决递归的好方式, Hlo模块的实现也比较简单, 可以认为是kset-kobj的结构的C++版本:在HloComputationProto里记录HloInstructionProto,在HloInstructionProto里记录被调用的另一个HloCProto实例Id

Continue reading

Tensorflow XlaOpKernel | tf2xla 机制详解

compiler/aot/ 以AOT的方式将tf2xla/接入TF引擎
compiler/jit/以JIT的方式将tf2xla/接入TF引擎, 核心是9个优化器和3个tfop,其中XlaCompileOp调用tf2xla的“编译”入口完成功能封装,XlaRunOp调用xla/client完成“运行”功能。
compiler/tf2xla/对上提供xla_compiler.cc:XlaCompiler::CompileFunction()供jit:compile_fn()使用将cluster转化为XlaComputation。核心是利用xla/client提供的接口,实现对XlaOpKernel的“Symbolic Execution”功能。每个XlaOpKernel子类均做的以下工作: 从XlaOpKernelContext中取出XlaExpression或XlaOp, 调用xla/client/xla_buidler.h提供的方法完成计算, 将计算结果的XlaOp存入XlaKernelContext.
compiler/xla/client/ 对上提供xla_builder.cc:Builder等供CompileFunction()使用,将Graph由Op表达转化为HloModuleProto:HloComputationProto:HloInstructionProto表达并保存在XlaComputation中。
对上提供local_client.cc:LocalClient::Compile(),作为编译入口供jit:BuildExecutable()使用,将已经得到的XlaComputation交给service并进一步编译为二进制。
对上提供local_client.cc:LocalExecutable::Run(),作为运行入口供jit/kernels/xla_ops.cc:XlaRunOp使用,通过Key找到相应的二进制交给service层处理
compiler/xla/service/ 对上提供local_service.cc:LocalService::BuildExecutable()供LocalClient::Compile()使用实现真正的编译,承接XlaComputation封装的HloProto, 将其转化为HloModule:HloComputation:HloInstruction表达, 对其进行优化之后, 使用LLVM后端将其编译为相应Executable后端的二进制代码
对上提供executable.cc:Executable::ExecuteOnStream()供LocalExecutable::Run()使用实现真正的执行二进制。

从Kernel的视角, XLA并不会新增Op, 而是针对已有的Op, 新增了基于XLA的另一个版本的Kernel: XlaOpKerne。在TF引擎中, OpKernel在软件栈上已是底层, 即最终的计算过程都要在OpKernel中实现. 但在XLA中, XlaOpKernel只是编译的入口, 大量的实际工作都交给了更下层的XLA引擎去完成.XLA相关的代码在tensorflow/compiler中.

tf2xla/负责XlaOpKernel的构造, 注册. 虽然XLA与TF引擎不在一层, 但二者面临的问题有很多有相似之处, 比如都需要对Kernel和Device保持易扩展性, 都需要维持前驱/后继Kernel的数据流和控制流关系. 基于类似的种种原因, XLA内部实现的注册XlaOpKernel的接口与TF引擎中注册OpKernel的风格十分相似, 同时, 其内部实现又有本质的不同, 而这些”不同”, 正是我们需要关注的.

要理解XlaOpKernel与OpKernel的不同, 关键在于了解”Symbolic Execution“.
先来看TF引擎, 它的OpKernel::Compute()方法要: OpKernelContext.Input()取输入数据 ==> 计算 ==> OpKernelContext.SetOutput()存输出数据, 计算结果继续通过OpKernelContext流入后继Opkernel, 其中流动的是真正的训练数据, 暂且将这个过程称之为”Execution”.
对比之下, XLA中的”Symbolic Execution”中的”Symbolic”即是说, XlaOpKernel的设计目的不在于去处理训练数据, 而在于去生成能够正确的处理数据的代码. 整个JIT类似于python解释器,先将程序编译为二进制,再运行二进制,XlaOpKernel执行的”Symbolic Execution”就类似其中的”编译为二进制”的过程. 具体地, 在XlaOpKernel::Compile()中: XlaOpKernelContext.Input()以XlaOp形式取输入 ==> 调用xla/client/xla_buidler.h提供的方法实现Op该有的功能, 实际上是生成一组能处理数据的HloInstruction ==> XlaOpKernelContext.SetOutput()存储XlaOp形式的结果, 计算结果继续通过XlaOpKernelContext流入后继XlaOpkernel, 其中流动的都是以XlaOp表征的对训练数据的处理方法.

//compiler/xla/client/xla_builder.h
// This represents an instruction that has been enqueued using the XlaBuilder.
// This is used to pass to subsequent computations that depends upon the
// instruction as an operand.
class XlaOp {

至于真正处理数据的时机, 就要交给XLA引擎, 它来负责后续的”编译”和”执行”, 具体地, 在JIT中, XlaCompileOp会在所有的XlaOpKernel::Compile()执行完毕之后, 继续调用xla/service中相应的方法将这些所有生成的HloInstruction编译生成二进制并进一步交给XlaRunOp来执行.

Continue reading

Tensorflow JIT 技术详解

compiler/aot/ 以AOT的方式将tf2xla/接入TF引擎
compiler/jit/以JIT的方式将tf2xla/接入TF引擎, 核心是9个优化器和3个tfop,其中XlaCompileOp调用tf2xla的“编译”入口完成功能封装,XlaRunOp调用xla/client完成“运行”功能。
compiler/tf2xla/对上提供xla_compiler.cc:XlaCompiler::CompileFunction()供jit:compile_fn()使用将cluster转化为XlaComputation。核心是利用xla/client提供的接口,实现对XlaOpKernel的“Symbolic Execution”功能。每个XlaOpKernel子类均做的以下工作: **从XlaOpKernelContext中取出XlaExpression或XlaOp, 调用xla/client/xla_buidler.h提供的方法完成计算, 将计算结果的XlaOp存入XlaKernelContext.**
compiler/xla/client/ 对上提供xla_builder.cc:Builder等供CompileFunction()使用,将Graph由Op表达转化为HloModuleProto:HloComputationProto:HloInstructionProto表达并保存在XlaComputation中。
对上提供local_client.cc:LocalClient::Compile(),作为编译入口供jit:BuildExecutable()使用,将已经得到的XlaComputation交给service并进一步编译为二进制。
对上提供local_client.cc:LocalExecutable::Run(),作为运行入口供jit/kernels/xla_ops.cc:XlaRunOp使用,通过Key找到相应的二进制交给service层处理
compiler/xla/service/ 对上提供local_service.cc:LocalService::BuildExecutable()供LocalClient::Compile()使用实现真正的编译,承接XlaComputation封装的HloProto, 将其转化为HloModule:HloComputation:HloInstruction表达, 对其进行优化之后, 使用LLVM后端将其编译为相应Executable后端的二进制代码
对上提供executable.cc:Executable::ExecuteOnStream()供LocalExecutable::Run()使用实现真正的执行二进制。

XLA基于编译技术将静态子图转换为二进制进而实现在某些网络下的加速,以BERT为例,P40单卡每batch性能从850ms提升到了700ms

JIT 是目前TF中两种XLA应用方式之一, 借助TF对计算图的先优化再执行的机制, JIT使用9个优化器+3个基于XLA模块实现的Tfop将Just In Time技术接入TF图计算引擎. 

==> Orignal Graph == Optimization ==> Optimized Graph ==> Graph Executor
                          |                                     |
                    JIT Add 7 Optimizations              XlaCompileOp and XlaRunOp
                    which Replace original               work just like other tfops
                    ops with 
                    (XlaCompileOp +XlaRunOp)s

9个Optimization如下图所示:

其中, MarkForCompilationPass, EncapsulateSubgraphsPass 和 BuildXlaOpsPass 最为关键.

//jit_compilation_pass_registration.cc, functionalize_control_flow_pass_registration.cc
REGISTER_OPTIMIZATION(OptimizationPassRegistry::PRE_PLACEMENT, 26, EncapsulateXlaComputationsPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::PRE_PLACEMENT, 25, IntroduceFloatingPointJitterPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::PRE_PLACEMENT, 27, FunctionalizeControlFlowPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, 5, CloneConstantsForBetterClusteringPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, 10, MarkForCompilationPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, 20, IncreaseDynamismForAutoJitPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, 30, PartiallyDeclusterPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, 40, EncapsulateSubgraphsPass);
REGISTER_OPTIMIZATION(OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, 50, BuildXlaOpsPass);

3个Op如下, 经过之前的优化, 已经完成了”图->XlaCompileOp + XlaRunOp”的转化, 而JIT的编译过程就在替换了原始Op的XlaCompile中进行. XlaCompileOp编译得到的二进制直接送到紧接其后的XlaRunOp中执行, 由于, XlaCompileOp里有用于存储之前编译结果的Cache, 所以理想情况下(图不变,输入的shape也不变), 只有第一次会真正的编译, 之后的step中由于Cache hit, XlaCompileOp的成本就很低了, 这也是XLA你能够实现加速核心原因. 据此, 在特征识别等输入频繁变动的场景, 由于XlaCompileOp的Cache Miss的概率大大增加, 整体性能就会比常规的TF执行引擎差. 

//jit/kernels/xla_ops.h
class XlaCompileOp : public OpKernel {
class XlaRunOp : public OpKernel {
class XlaLocalLaunchOp : public XlaLocalLaunchBase : public OpKernel { //for eager

Optimization

Continue reading

Tensorflow JIT/XLA UML

compiler/aot/ 以AOT的方式将tf2xla/接入TF引擎
compiler/jit/以JIT的方式将tf2xla/接入TF引擎, 核心是9个优化器和3个tfop,其中XlaCompileOp调用tf2xla的“编译”入口完成功能封装,XlaRunOp调用xla/client完成“运行”功能。
compiler/tf2xla/对上提供xla_compiler.cc:XlaCompiler::CompileFunction()供jit:compile_fn()使用将cluster转化为XlaComputation。核心是利用xla/client提供的接口,实现对XlaOpKernel的“Symbolic Execution”功能。每个XlaOpKernel子类均做的以下工作: **从XlaOpKernelContext中取出XlaExpression或XlaOp, 调用xla/client/xla_buidler.h提供的方法完成计算, 将计算结果的XlaOp存入XlaKernelContext.**
compiler/xla/client/ 对上提供xla_builder.cc:Builder等供CompileFunction()使用,将Graph由Op表达转化为HloModuleProto:HloComputationProto:HloInstructionProto表达并保存在XlaComputation中。
对上提供local_client.cc:LocalClient::Compile(),作为编译入口供jit:BuildExecutable()使用,将已经得到的XlaComputation交给service并进一步编译为二进制。
对上提供local_client.cc:LocalExecutable::Run(),作为运行入口供jit/kernels/xla_ops.cc:XlaRunOp使用,通过Key找到相应的二进制交给service层处理
compiler/xla/service/ 对上提供local_service.cc:LocalService::BuildExecutable()供LocalClient::Compile()使用实现真正的编译,承接XlaComputation封装的HloProto, 将其转化为HloModule:HloComputation:HloInstruction表达, 对其进行优化之后, 使用LLVM后端将其编译为相应Executable后端的二进制代码
对上提供executable.cc:Executable::ExecuteOnStream()供LocalExecutable::Run()使用实现真正的执行二进制。




Related:
Tensorflow XLA Service Buffer优化详解
Tensorflow XLA Service 详解 II
Tensorflow XLA Service 详解 I
Tensorflow XLA Client | HloModuleProto 详解
Tensorflow XlaOpKernel | tf2xla 机制详解
Tensorflow JIT 技术详解
Tensorflow JIT/XLA UML
Tensorflow OpKernel机制详解
Tensorflow Op机制详解
Tensorflow Optimization机制详解
Tensorflow 图计算引擎概述

Tensorflow XLA Debug/Profiling Methods

Tensorflow 调试方式主要有以下几种

  1. Timeline初步分析性能瓶颈
  2. Tensorboard,Graphviz等对图信息进行可视化
  3. 环境变量配置导出图信息
  4. 使用bazel编译选项输出内部信息
  5. 使用GDB进行细节问题定位

Timeline

可以直观的显示一个Step内Op粒度的执行时间, 为性能优化提供参考。Chrome中键入 chrome://tracing/ ,即可通过Load按钮加载time.json, 使之图形化显示

python tf_cnn_benchmarks.py --trace_file="time.json"  --save_summaries_steps=10 --summary_verbosity=1 --train_dir "./tmp"  --num_gpus=1 --batch_size=1 --model=alexnet
mpirun --allow-run-as-root -np 4 -H 10.54.32.198:4 python tf_cnn_benchmarks.py --xla="True" --trace_file="time.json" --num_gpus=1 --batch_size=32 --model=alexnet --variable_update=horovod

Timeline看起来很美好,但有几个问题一定要注意:

  1. 基本单位是Op。 Op内逻辑再简单,也一定会体现,Op内逻辑再复杂,也不会对其拆分。
  2. 是谁的执行流。Timeline体现的是CPU端启动Op的执行流,,上图中的GPU端的执行流看起来有多条,实际上到底有没有并行呢?这不一定,这张Timeline图本质只是对time.json进行了可视化,这里看到的GPU任务就对应time.json的 "name": "/job:localhost/replica:0/task:0/device:GPU:0 Compute" 而GPU的执行流编号就对应 "tid": 14, 所以,问题就变成了:这里的“tid”能否表示实际在GPU中的执行流?我理解是不能的,在使用GPU做backend的场景,HorovodAllreduce底层调用NcclAllReduce,而NcclAllReduce本质是DMA引擎在做实际工作,而以DMA的工作方式,怎么会占用GPU的Compute流呢,所以,这里应该比较准确的理解应该是,TODO:在某一时刻,不断开启新的执行流去执行一个GPU Op,至于这些执行流是否真的在GPU内并行,这里并不能表达,需要使用GPU专属的工具。
  3. 并行的欺骗性,本质上,这种误导是由于1和2导致的Timeline局限性,倒也称不上是错误。上图中使用Horovod进行分布式训练时,由于多个AllreduceOp并无相关性,也分属不同的执行流,看起来并行度很高,其实在其内部实现中,一方面Horovod会做融合,另一方面每个HorovodAllreduceOp只有一个后台线程真正的去做网络通信,最后,网络IO的能力上限取决于网卡,当网卡达到瓶颈时,更多的并行HorovodAllreduceOp是没有意义的。
  4. 和Step强相关,取一个Timeline的时候,通常会指定运行到某个Step来保存,这里的问题是,不同的Step可能执行的Timeline差异巨大,典型的问题就是开启XLA时,部分子图只会在Step1执行,同时,Step1进行的XlaCompile也会大大延长单Step时间,所以在测试的时候,使用的Step也是要小心设置的。

Tensorboard

Tensorboard用于传统TF Graph调试是比较容易理解的,这里讨论如何使用Tensorboard进行XLA的可视化。XLA可以将优化的各个阶段的子图导出为pbtxt或pb|dot,html等格式,如果是pbtxt,可以直接点击Tensorboard左侧Choose File 按钮上传。

python tf_cnn_benchmarks.py --trace_file="time.json" --xla="True" --save_summaries_steps=10 --summary_verbosity=1 --train_dir "./tmp"  --num_gpus=1 --batch_size=1 --model=alexnet

执行”./bazel-bin/tensorboard/tensorboard –logdir=”./logs” –port=9000″,浏览器输入”http://localhost:9000/“,点击上传mark_for_compilation_annotated_1.pbtxt即可看到子图:

Dot + Graphvis

XLA的HLO PASS可以生成pb, html,dot等格式,这里的pb用import_pb_to_tensorboard或者自己写的脚本都没能成功转化成pbtxt,还没办法使用的Tensorboard查看,但相应的dot格式,显然可以使用Graphviz进行可视化。但鉴于graphviz的全量计算的方式,对于特别大神经网络,dot工具本身就会成为瓶颈。

foo $ dot -Tsvg -o module_0002.after_optimizations.svg module_0002.after_optimizations.dot
Continue reading

Tensorflow 编译方法

Tensorflow使用SWIG自动根据so文件生成相应的python接口, 这就引起一系列反常识的问题, 比如, TF的编译环境同时要求TF部分可运行, 通常来说, 编译环境只需要配置编译依赖即可, 但在TF中行不通, 举个例子: 如果环境中只有CUDA, 没有GPU以及可用的Driver, 那么在编译的最后阶段, SWIG机制会去运行Tensorflow Runtime来生成Python接口, 但缺失GPU或Driver的环境是无法运行TF的, 也就无法完成位于编译最后阶段的python接口生成. 此外, 这也导致了如果在SWIG调用路径上的Runtime代码一旦有逻辑问题而不是语法问题, 在Runtime的编译过程不会报错, 编译最后生成Python接口也会出错导致最终编译失败, 进一步导致该类逻辑错误难以调试.

Debug信息保留

将编译命令调整如下

bazel --output_user_root="/.cache/bazel/" build --config=cuda --config=nccl --copt=-g  --strip=never //tensorflow/tools/pip_package:build_pip_package

第一处`–output_user_root=”/.cache/bazel/”` 主要解决编译环境root一般较小, 无法承载编译临时文件
第二处` –copt=-g  –strip=never ` 用于保留调试信息, 要配合python3-dev使用, 如果使用python2, 由于带有debug信息生成so过大, 会导致wheel制作失败(int32溢出), 当然, 相应的python代码也要适配python3进行修改

离线安装依赖

Tensorflow编译过程会从Internet上load所需依赖, 如果网络环境不允许, 可以通过其他路径现将依赖包下载copy到编译环境, 再修改tensorflow/tensorflow/workspace.bzl中的相关url, 即可规避联网编译问题

https://github.com/tensorflow/tensorflow/issues/5428#issuecomment-258732841

tf_http_archive(
    name = "eigen_archive",
    build_file = clean_dep("//third_party:eigen.BUILD"),
    patch_file = clean_dep("//third_party/eigen3:gpu_packet_math.patch"),
    sha256 = "0dde8fb87f5dad2e409c9f4ea1bebc54e694cf4f3b633081b0d51a55c00f9c9f",
    strip_prefix = "eigen-eigen-a0d250e79c79",
    urls = [
        "http://mirror.tensorflow.org/bitbucket.org/eigen/eigen/get/a0d250e79c79.tar.gz",
        #"https://bitbucket.org/eigen/eigen/get/a0d250e79c79.tar.gz",
        "file:////tmp/sources/tensorflow/a0d250e79c79.tar.gz",
    ],
)

Tensorflow OpKernel机制详解

OpKernel是Op的具体实现, tf中已经实现的tfop的OpKernel在源码中的tensorflow/core/framework/kernel/中, OpKernel通过注册时使用相同的名字将自己和相应的Op联系到一起.

在tf中, OpKernel进一步可以分为两类, OpKernelAsyncOpKernel:

  1. OpKernel是同步执行的, 即"Compute()"返回即认为数据已经被正确处理, 注册OpKernel, 子类需要重写其Compute()方法.
  2. AsyncOpKernel是对OpKernel的封装, 顾名思义, AsyncOpKernel执行返回并不意味着数据已经被处理完毕, 数据的真正被处理完毕时通过回调的方式通知Op执行引擎, 注册一个AsyncOpKernel, 子类需要实现"AsyncCompute()"而不是Compute().

接口形式

无论是哪种OpKernel, 均使用"REGISTER_KERNEL_BUILDER()"注册到运行核心.

//tensorflow/core/common_runtime/kernels/nccl_ops.cc
#include "third_party/nccl/nccl.h"
#include "tensorflow/core/framework/op_kernel.h"
#include "tensorflow/core/nccl/nccl_manager.h"
namespace tensorflow {
class NcclAllReduceOpKernel : public AsyncOpKernel {
 public:
  explicit NcclAllReduceOpKernel(OpKernelConstruction* c)
      : NcclReduceOpBase(c) {}

  void ComputeAsync(OpKernelContext* c, DoneCallback done) override {
    //...
  }
};
REGISTER_KERNEL_BUILDER(Name("NcclAllReduce").Device(DEVICE_GPU),
                        NcclAllReduceOpKernel);  
}

注册原理

注册机制的实现代码主要集中在tensorflow/core/framework/op_kernel.h(.cc). 与 Optimization以及Op在注册时直接构造一个static OptimizationPassRegistration(OpRegistrationData)对象的机制略有不同, OpKernel的通过一些trick实现了对OpKernel的延迟构造, 即"REGISTER_OP_KERNEL_BUILDER()"并没有直接构造一个"OpKernel"实例, 而是构造一个" static ::tensorflow::kernel_factory::OpKernelRegistrar "对象, 并借由该构造过程构造并注册一个"KernelRegistration" 对象到 global_regsitry, 该构造过程接受上层传入的, 用于new一个OpKernel的"[](::tensorflow::OpKernelConstruction* context) -> ::tensorflow::OpKernel* { return new __VA_ARGS__(context);}" 函数, 在上层真正需要这个OpKernel的时候, 才会通过一系列调用最终执行该"create_fn()/lambda"来构造一个实实在在的OpKernel对象. 即整体上不再是Registry->Registration(Optimization/Op对象), 而是 Registry->Registrar->Registration->在需要时create_fn()构造OpKernel对象.

构造一个OpKernelRegistrar:

Continue reading

Tensorflow Op机制详解

注册一个tfop分为两部分: OpOpKernel. 其中, Op是tfop的声明部分, 类似于函数的声明, 主要描述Op静态属性. OpKernel是tfop的实现部分, 同样类似于函数的实现, 主要描述OpKernel的具体计算逻辑. tf已经实现的Op相关源码在tensorflow/core/common_runtime/ops/中, OpKernel的部分在tensorflow/core/common_runtime/kernels/. 二者通过注册时使用的名字联系到一起.

接口形式

下面是一个注册Op的例子, 使用接口"REGISTER_OP()"‘注册一个Op,包括输入输出, Op属性等内容, 注册的Op最终会转换为OpDef, 即用protobuf格式存储的静态数据, 所以, 这里用到的属性也按照pb的格式以KV的形式编写. REGISTER_OP的实现在”tensorflow/core/framework/op.h OpDefBuilderWrapper” 中, 众多可设值中,"Attr"是一个允许自定义的值, 比如XLA引擎就根据自身需求提供了"XlaCompile", 如果一个Op将该值设置为true, 就会强制XLA引擎将其编译. 当然, 也可以设置一些无用的值, 就像函数声明里有一个并没有实际使用的参数, 除了浪费存储空间没有其他用途.

//tensorflow/core/common_runtime/nccl_ops.h
#include "tensorflow/core/framework/common_shape_fns.h"
#include "tensorflow/core/framework/op.h"

namespace tensorflow {

using shape_inference::InferenceContext;
using shape_inference::ShapeHandle;

REGISTER_OP("NcclAllReduce")
    .Input("input: T")
    .Output("data: T")
    .Attr("reduction: {'min', 'max', 'prod', 'sum'}")
    .Attr("T: {half, float, float64, int32, int64}")
    .Attr("num_devices: int")
    .Attr("shared_name: string")
    .Attr("XlaCompile: bool=true")
    .SetIsStateful()
    .SetShapeFn(shape_inference::UnchangedShape);
}

注册原理

在内部原理上, 上述代码可以进一步分解为两部分: 构造一个Op + 注册一个Op, 这部分代码主要在 “tensorflow/core/framework/ op.h(.cc) op_def_builder.h(.cc) ” 中

Continue reading

Tensorflow OptimizationPassRegistry机制详解

整个流水线大概是:Init Graph –>OptimizationPassRegistry和Grappler进行全图优化 –> 根据Device将Graph拆成 Partition –> GraphOptimizer优化Partition–>图执行. 同其他注册机制一样, OptimizationPassRegistry也使用的registry, registertion以及registerar等概念.

源码中, 使用REGISTER_OPTIMIZATION()注册一个优化器, 具体实现如下

//core/common_runtime/optimization_registry.h
#define REGISTER_OPTIMIZATION(grouping, phase, optimization) \
  REGISTER_OPTIMIZATION_UNIQ_HELPER(__COUNTER__, grouping, phase, optimization)

#define REGISTER_OPTIMIZATION_UNIQ_HELPER(ctr, grouping, phase, optimization) \
  REGISTER_OPTIMIZATION_UNIQ(ctr, grouping, phase, optimization)

#define REGISTER_OPTIMIZATION_UNIQ(ctr, grouping, phase, optimization)         \
  static ::tensorflow::optimization_registration::OptimizationPassRegistration \
      register_optimization_##ctr(                                             \
          grouping, phase,                                                     \
          ::std::unique_ptr<::tensorflow::GraphOptimizationPass>(              \
              new optimization()),                                             \
          #optimization)

-12-new了一个我们注册的optimization对象并用unique_ptr指向它, 这个unique_ptr就是registry管理的对象, 通过它间接管理相应的optimization. 注册的本质是返回一个静态的, 类型为'OptimazationPassRegistration'的, 名为'register_optimization_##ctr'的对象, 这里使用了C++预编译宏'__COUNTER__'生成唯一变量名

下面是一个`OptimazationPassRegistration`对象的构造过程, 可以看出, 就是将我们构造的'register_optimization_##ctr注册到全局的'global_optimization_registry'.

Continue reading

Tensorflow 图计算引擎概述

tensorflow的用户可以使用多种语言来构造的自己的图, 但各种语言的API最终都会经由C API 进入tensorflow 运行时. 可以说, 对于运行时代码, 其上边界就是C API. 比如, 通过python描述的一张网络, 就是通过类似下面的几个python-C接口进入运行时的.

#9  0x00007f9de118daa4 in PyEval_EvalFrameEx ()
#10 0x00007f9de118f0bd in PyEval_EvalCodeEx ()

tensorflow整体上可以看做一个”图语言的编译器”, 和所有编译器的优化以及翻译的功能类似, Graph在运行时中的处理过程, 可以分为 图构造->图优化->图执行 几个阶段. 其中, 图优化随同图构造一同被执行.

全图构造及其优化

Session初次构造时, 应用层代码中定义的数据流图转换成GraphDef格式, 经由C API传入DirectSession.Extend(), 参考调用栈如下

PyEval_EvalCodeEx()
  PyEval_EvalFrameEx()
    _wrap_ExtendSession()
      tensorflow::ExtendSession()
        tensorflow::ExtendSessionGraphHelper()
          tensorflow::SessionRef::Extend()
            tensorflow::DirectSession::ExtendLocked()
              tensorflow::DirectSession::MaybeInitializeExecutionState(out_already_initialized/already_initialized)
                if out_already_initialized:
                  return
                  flib_def_.reset(new FunctionLibraryDefinition())
                tensorflow::GraphExecutionState::MakeForBaseGraph()
                    std::unique_ptr<GraphExecutionState> ret(new GraphExecutionState(graph_def, options));
                    if (!ret->session_options_->config.graph_options().place_pruned_graph()):
                      tensorflow::GraphExecutionState::InitBaseGraph()
                        OptimizationPassRegistry::Global()->RunGrouping(PRE_PLACEMENT)
                        Placer placer()
                        placer.Run()
                        OptimizationPassRegistry::Global()->RunGrouping(POST_PLACEMENT)
                out_already_initialized = false
              if already_initialized:
                flib_def_->AddLibrary(graph.library())
                std::unique_ptr<GraphExecutionState> state
                execution_state_->Extend(graph, &state))
                execution_state_.swap(state)
    _wrap_TF_SessionRun_wrapper()
      tensorflow::TF_SessionRun_wrapper()
        tensorflow::TF_SessionRun_wrapper_helper()
          TF_SessionRun()
            TF_Run_Helper()
              tensorflow::SessionRef::Run()
                tensorflow::DirectSession::Run()  
                  DirectSession::GetOrCreateExecutors(executors_and_keys)
                    CreateExecutors()
                      std::unique_ptr<ExecutorsAndKeys> ek(new ExecutorsAndKeys); 
                      std::unordered_map<string, std::unique_ptr<Graph>> graphs;
                      CreateGraphs(&graphs)
                        if (options_.config.graph_options().place_pruned_graph()):
                          MakeForPrunedGraph()
                            ret->InitBaseGraph()
                              if (session_options_ && session_options_->config.graph_options().place_pruned_graph()):
                                PruneGraph()
                                  if (options.use_function_convention):
                                    feed_rewrites.emplace_back(new subgraph::ArgFeedRewrite())
                                    fetch_rewrites.emplace_back(new subgraph::RetvalFetchRewrite())
                                    ValidateFeedAndFetchDevices()
                                  else:
                                    feed_rewrites.emplace_back(new subgraph::RecvFeedRewrite())
                                    fetch_rewrites.emplace_back(new subgraph::SendFetchRewrite())
                                  subgraph::RewriteGraphForExecution(graph, feed_rewrites, fetch_rewrites)
                              OptimizationPassRegistry::Global()->RunGrouping(PRE_PLACEMENT)
                              Placer placer()
                              placer.run()
                              OptimizationPassRegistry::Global()->RunGrouping(POST_PLACEMENT)
                            graph_ = new_graph.release();
                            ret->BuildGraph()
                              OptimizeGraph()
                                grappler::RunMetaOptimizer()
                                  MetaOptimizer::Optimize()
                              if (session_options_ == nullptr || !session_options_->config.graph_options().place_pruned_graph()) {
                                PruneGraph()
                              std::unique_ptr<ClientGraph> dense_copy(new ClientGraph)
                        else:
                          execution_state->BuildGraph();
Continue reading

BERT 模型参数量估计

根据BERT论文, 其12层transformer结构有110M参数, 24层更是高达340M, 虽然google公开了这两个网络的预训练模型, 用户只需在后面加一层Full-connected 但如果自己去做BERT预训练, 到底要耗费多少显存呢?

tensorflow中使用tf.variable()在模型中生成训练参数, 同时, tf.dense()内部也含有weight和bias两类参数. 比如,

tf.get_variable(shape=[vocab_size, embedding_size],initializer=…))

生成了shape是[vocab_size, embedding_size]的variable, 即 vocab_size * embedding_size 个参数,

to_tensor_2d = tf.layers.dense(from_tensor_2d, kernel_initializer=...)) 

内含 from_tensor_2d.shape[-1] * to_tensor_2d.shape[-1] 个weight 以及 to_tensor_2d.shape[-1] 个bias. 共计 (from_tensor_2d.shape[-1] + 1) * to_tensor_2d.shape[-1] 个参数, 这就是一个FC-Layer的参数量

有了上面的例子, BERT的参数量就不难分析, 整个BERT网络可以分为4个模块: embedding layer + transformer layers + pooled layer+ classifier layer

embedding-layer

Continue reading

tfrecord二进制解析

tfrecord是tensorflow基于protobuf框架开发的一种用于持久化训练数据的文件。

protobuf

考虑下面一个简单的protobuf message定义:

syntax = "proto3";
message Msg{
	map<string, int64> val = 1;
};

这里,我们将val的field_number属性置为1,下文将在序列化文件中将其提取出来。通过 protoc --proto_path=. --python_out=. demo.proto编译生成demo_pb2.py,然后就可以使用python脚本按照Msg格式对数据进行编解码。这里,我们使用的序列化代码如下:

import demo_pb2
import sys
v = {
    "test":300,
}
msg= demo_pb2.Msg(val = v)
fo=open(sys.argv[1], "wb")
fo.write(msg.SerializeToString())
fo.close()

该代码的就是将数据v持久化到传入的文件中,其中encode接口SerializeToString()是在编译处demo_pb2.py自动生成的,与之相对应的还有decode接口ParseFromString()。 在python write.py data之后,我们可以解析data的二进制格式:

# xxd -b data
0000000: 00001010 00001001 00001010 00000100 01110100 01100101  ....te
0000006: 01110011 01110100 00010000 10101100 00000010           st...

在pb协议中,任何数据都是由key-value的形式管理,其中,key即message中每个字段的field_number, 除了key,还要有wire_type属性, field_number<<3|wire_type构成一个字节,对于有wire_type == length-delimited的数据,还需要一个字节描述数据长度的payload属性。 我们逐个字节解释这几行二进制:

[00001010]Msg's key,经过逆运算就可以获取到field_number==1,wire_type=2(Length-delimited)
[00001001]Msg's payload:9 byte
[00001010, 00000010(end)]Msg's value。
{
    {
        [00001010]Msg.map<>.string's key,可知其field_number == 1, wire_type == 2
        [00000100]Msg.map<>.string's payload:4 byte
        [01110100 01100101 01110011 01110100]Msg.map<>.string's value:test
    }
    {
        [00010000]Msg.map<>.int64's key:field_number == 2, wire_type == 0(varint)
        [10101100, 00000010(end)]Msg.map<>.int64's value:300
    }
}
Continue reading

Linux 任务调度简介

吞吐 和 响应 是评估一个计算机系统性能的主要参数,比较典型的就是在IO系统中,Linux为了提高IO吞吐而将IO先缓存到内存中,这在本质上是牺牲了对某个IO的响应过程并以此来提高吞吐率。这种设计思路在Linux调度器的设计也可以看到,Linux并没有追求硬实时系统的高响应,而是在吞吐和响应之间取得一个平衡。

在Linux中,任务(或者说task_struct,由于Linux中进程和线程没有明显的界限,本文使用任务来描述一个可调度实例)一共有0-139个优先级,其中0-99是给RT 任务使用,100-140给普通任务使用,前者即实时调度策略,进一步又分为SCHED_FIFO和SCHED_RR,后者即普通任务,即SCHED_NORMAL。作为优先级更高的实时任务,SCHED_FIFO和SCHED_RR的任务一旦处于RUNNING(就绪)状态,就可以无条件抢占SCHED_NORMAL的CPU,而只要系统中有RT的任务在运行,SCHED_NORMAL就没有运行的机会。当然,即便是RT 任务,其也会有优先级高低的区别,对于不同优先级的RT任务,显然要严格按照高优先级优先运行的原则,而对于相同优先级的两个RT任务,如果使用SCHED_FIFO,就按照”先来后到”的顺序,只要先运行的任务没有放弃CPU,后来的任务不会被调度到; 如果使用SCHED_RR,同优先级的任务会”周期性”的都被执行到。

Linux中每一个SCHED_NORMAL任务都有一个nice值(-20~+19,缺省为0),这个nice值表示这个任务对于其他任务的”容忍”程度,nice值越大,表示这个任务越好说话,结果是这个任务占用的资源就会减少。早期的Linux通过动态的调整一个动态nice来平衡IO消耗型任务和CPU消耗型任务,不过如今这个工作已经被巧妙的CFS替代了

vruntime = pruntime/weight * delta

上图就是CFS的运作公式,pruntime表示当前任务实际运行的时间,weight是一个与nice值相关的权重,share是系统级/cgroup的调整参数。整个CFS基于红黑树实现进程调度。如下图所示:

上图的数字就是每个任务的vruntime,CFS在进行任务调度的时候,总是调度红黑树中vruntime最小的那个任务,由于IO消耗型的任务本身pruntime就比CPU消耗型的pruntime小,所以整体上IO消耗型的任务的vruntime更小,也就能够获得更多的被调度的机会,这样就能提高系统整体的吞吐。

Continue reading

Linux 内存管理 VI

内存统计是资源评估的重要方面,本文从进程和系统两个视角讨论smem和free两个常用的内存统计工具,Linux进程虚拟地址空间和物理内存的关系如下。

进程视角

进程内存消耗通常使用VSS、RSS、PSS、USS来表征,需要注意的是,进程的这些指标的统计是不包括内核地址空间的。
VSS = 1 + 2 + 3 virtual set size,进程的虚拟地址空间都是通过 task_struct->mm_struct ->mmap->vm_area_struct 来管理其虚拟地址空间,每一个VMA都对应一个Section,所有的VMA对应到进程的内存消5就是VSS,以下两点需要注意:

  1. VSS不是真的体现进程物理内存的消耗,而是虚拟地址空间的消耗。
  2. 对于共享库来说,代码段只有一份,数据段是数据调用进程的,每个进程一份

可以通过如下方式查看VSS(VSZ字段)

$ps aux|more
USER       PID %CPU %MEM    VSZ   RSS TTY      STAT START   TIME COMMAND
root         1  0.0  0.0 185552  5512 ?        Ss   2月16   0:01 /sbin/init noprompt persistent splash
root         2  0.0  0.0      0     0 ?        S    2月16   0:00 [kthreadd]
root         3  0.0  0.0      0     0 ?        S    2月16   0:01 [ksoftirqd/0]

RSS = 4 + 5 + 6,resident set size,只进程实际占用的内存,对于Lib,其所占内存完整的体现在每一个调用的进程RSS中
PSS = 4/3 + 5/2 + 6, proportional set size,同样是体现进程实际占用的内存,不同是对于多进程公共的Lib,按照比例将其占用内存摊派到每一个进程的PSS中
USS = 6 unique set size ,进程独占的内存,比如进程的Heap,分析内存泄漏时通常使用USS。
可以通过如下方式查看USS、PSS、RSS,smem工具通过读取**/proc//smaps和maps**两个文件整理输入如下。

$smem -P ^chromium
  PID User     Command                         Swap      USS      PSS      RSS 
16317 jiang    /usr/lib/chromium-browser/c        0      164     2022    13764
Continue reading