TVM: An Automated End-to-End Optimizing Compiler for Deep Learning

Woojin LEE

2021.02.21 | Research, KAIST EE Co-op, Internship

full

0. Disclaimer

Apache TVM은 현재도 꾸준히 개발중인 프로젝트이기에 홈페이지, 코드 등이 빠르게 바뀌고 있습니다. 본 포스트는 2020/08 ~ 2021/01에 작성된 글입니다. 최신의 내용과 링크/내용 등이 달라질 수 있음을 미리 밝힙니다.
본 포스트는 TVM을 처음 접하는 사람들에게 주안점 등을 제시하는 역할을 하는 것이 목표입니다. TVM을 처음 보시는 분들도 이 글을 통해 어느 정도 큰 그림을 그릴 수 있었으면 좋겠습니다. 자료의 일부는 논문에서, 일부는 홈페이지에서, 일부는 컨퍼런스 자료에서 발췌하였고, 이들을 바탕으로 한 번역+가이드라인 정도라고 보면 될 것 같습니다. TVM echosystem에 작은 도움이 될 수 있다면 행복할 것 같습니다😃

1. Introduction to TVM

최근 딥러닝에 관련된 새로운 기술들이 무서운 속도로 등장하고, 또 발전하고 있습니다. Tensorflow, Pytorch 등등 다양한 DL framework 뿐 아니라, 다양한 형태의 하드웨어들도 각자의 영역에서 꾸준히 발전해가고 있습니다.

NVIDIA CUDA처럼 하드웨어 벤더에서는 그 하드웨어를 잘 utilize하기 위한 라이브러리가 존재합니다. 이 하드웨어 라이브러리들은 Pytorch/Tensorflow 등에 있는 다양한 operator들을 효율적으로 작동시키기 위해 꾸준히 유지/보수해야 합니다. 하지만, CPU/GPU/FPGA/ASIC 등등 각각의 하드웨어는 서로 메모리 구조/데이터타입 등등 다양한 차이점들을 가지고 있기 때문에 하드웨어 라이브러리 최적화는 엄청난 비용을 요구합니다.

Wafer-scale-chip이라는, 혁신적인 칩(칩이 맞나...?)을 들고 나온 Cerebras라는 회사에서도 그 거대한 칩을 위한 컴파일러를 가지고 있습니다! (웨이퍼 크기의 칩을 잘 utilize할 수 있는 칩이 있다는 사실이 신기해서 TMI로 넣어보았습니다) 새로운 하드웨어를 개발하기 위해 기존의 Pytorch, Tensorflow 등과의 연동성을 담당하는 컴파일러까지 하드웨어 벤더에서 전부 개발하고 있었음을 보여주는 또 다른 하나의 예시가 될 수 있겠군요.

아래 영상에서 Theirry(VTA 개발자)의 박사학위 취득을 위한 ~~처절한~~ 이야기를 하며 쉽게 설명하였습니다.

이는 너무 비효율적이고, DL frameowrk과 하드웨어의 간극을 효과적으로 메꿀 수 있는 컴파일러에 대한 필요성이 점차 확대되고 있습니다.

Apache TVM은 이에 대한 솔루션으로, 딥러닝을 위한 Automated, End-to-End Optimizing OpenSourceCompiler입니다. 현재 AWS, 알리바바, 퀄컴, AMD 등 다양한 기업들에서도 TVM을 활용하는 사례가 늘고 있습니다.

ML2의 저널클럽에서 TVM에 관련하여 다룬 적이 있습니다. 그때 다루었던 슬라이드를 첨부합니다 :
🗂️ 슬라이드

2. TVM Workflow and Components

ninety

출처: https://tvm.apache.org/

TVM의 workflow와 주요 component를 알아보도록 하겠습니다. 자세한 내용은 TVM Docs-Design and Architecture와 각 섹션에 추가한 "🔎 함께 곁들이면 좋은 문헌"을 참조 바랍니다.

2.1. Overall Workflow

2.2절 Logical Components와 함께 보면 좋습니다.

ninety

Example Compilation Flow.(출처 : https://tvm.apache.org/docs/dev/index.html)

위 다이어그램을 바탕으로 TVM의 workflow를 간단히 살펴보자면,

DL framework(Tensorflow, Pytorch, Caffe 등)으로부터 모델을 import하여, relay를 통해 TVM의 IRModule(Intermediate Representation Module)로 변환합니다. IRModule에는 모델을 나타낼 수 있는 다양한 함수들이 포함되어 있습니다.
transformation과정으로, IRModule을 점점 lowering합니다. IRModule이 원하는 하드웨어에 알맞게(메모리 구조, 데이터 타입 등) 최적화될 수 있도록 Tensorize 시킵니다. 이를 위해 quantization, graph-packing 등 다양한 작업들이 포함되고, 많은 최적화 작업에 AutoTVM이 사용됩니다.
translation(codegen) 과정입니다. 지금까지는 전부 IRModule만 가공했습니다. 하지만, 우리의 궁극적인 목표는 타깃 하드웨어에서 잘 작동하는 코드입니다. 예를 들어 타깃 하드웨어가 GPU라면 [CUDA/OpenCL/Metal 등], CPU라면 [llvm] 코드를 생성해야 합니다. TVM에서 이는 runtime.Module에 encapsulate 되어있고, 이를 통해 타깃 하드웨어에 쉽게 export/load/execute시킬 수 있습니다.
runtime.Module를 로드하여 타깃 하드웨어에서 모델을 작동시킵니다.
Transformation

Transformation 과정은 크게 Optimization 과정과 Lowering 과정으로 나눌 수 있습니다. \ Optimization 과정은 같은 코드를 더 효율적으로 만드는 과정이고, Lowering 과정은 C언어로 작성된 파일이 어셈블리를 거쳐 기계어로 번역되듯, 타깃 하드웨어에 잘 맞게 돌아갈 수 있도록 하는 부분입니다.
- Optimization
  - relay/transform을 보면 이들은 모델을 최적화시키는 pass들을 정의하고 있습니다. 이는 일반적인 컴파일러들이 수행하는 최적화 과정(constant folding, dead-code elimination 등) 뿐 아니라 tensor 연산들도(layout transformation, scaling factor folding 등) 최적화시켜줍니다.
  - optimization 막바지 단계에는 큰 덩어리(end-to-end function, e.g. MobileNet)를 작은 덩어리들(sub-function, e.g. conv2d-relu)로 잘게 나누는, FuseOps라는 pass를 지나게 됩니다. 작게 분할한 뒤에는, 1)end-to-end function을 sub-function들의 시퀀스로 나타내는 과정과 2)각각의 sub-function 컴파일/최적화 과정을 거칩니다.
- Lowering
  - tir/transform는 tir에 대한 transformation pass를 담고 있습니다. 이 tir은 타깃 하드웨어에 맞게 multi-dimensional access를 one-dimensional pointer access로 바꾸거나 runtime calling convention에 맞추는 등 lowering을 목표로 합니다. (lowering만 있는 것은 아니고, optimization 과정들도 포함됩니다.)
사실, LLVM이나 CUDA C 등 기존 컴파일러들도 low-level optimization을 수행할 수 있습니다. TVM은 기존 컴파일러들이 잘 하는 register allocation 등의 문제는 이들에게 맡기고, 나머지 optimization들에 대해서만 다룬다고 합니다.
Translation

IRModule을 타깃 하드웨어에서 실행 가능하도록 만드는 부분입니다. 특정 하드웨어에 대해서는 external code generator를 통해 relay function의 sub-graph 부분을 직접적으로 translate할 수 있다고 합니다.(ex. x86이나 ARM을 사용하는 경우, LLVM IRBuilder를 사용하여 in-memory LLVM IR을 만듭니다.)
이 과정에서 [vector 길이, vectorization] 등등 하드웨어의 정보를 알려주기 위해 Target structure에 타깃의 정보를 담아 TVM에게 알려줍니다.

이 두 가지는 꼭 염두에 두고 흐름을 따라가시면 좋을 것 같습니다:

컴파일에서 주요 data structure 2가지
1. IRModule: contains relay.Function and tir.PrimFunc
2. runtime.Module: contains runtime.PackedFunc
IRModule과 runtime.Module간 transformation 2가지
1. relay/transform and tir/transform: determinstic, rule-based
2. auto_scheduler and autotvm: search-based

2.2. Logical Components

ninety

TVM Architecture Design (출처: <"https://tvm.apache.org/docs/dev/index.html">)

TVM을 구성하는 Logical Component 간 관계도입니다. 원문은 TVM docs에서 확인할 수 있습니다. 중요하다고 생각하는 몇 가지 부분만 나열하였습니다. 여러가지 컴포넌트들이 복잡하게 얽혀 이해하기 힘들 수 있습니다.

Logical component

⚙️ tvm/runtime

TVM supports multiple programming languages for the compiler stack development and deployment.

TVM이 지켜야 하는 사항에는 다음이 있습니다:

Deployment

Debug

Link

Prototype

Expose

Experiment

즉, 다양한 곳에서 서로를 호출하기 용이한 함수를 작성해야 하고, deploy할 디바이스에 들어갈 runtime core가 최대한 가벼워야 합니다.\ 이를 위해 tvm runtime이 작성되었습니다.

tvm.runtime.Module

컴파일 결과를 encapsulate하고, load/execute할 수 있습니다. 이름으로 호출하기 위해 GetFunction 메소드를 가지고 있습니다.

tvm.runtime.PackedFunc

type-erased function interface입니다. type-erased 되어있기 때문에 input이나 return 타입에 제한이 없습니다.

The restriction makes the implementation simple without the need of serialization

다만, TVM에서는 serialization 없이 구현할 수 있도록 하기 위해 PackedFunc에 대한 input과 output의 데이터 타입을 제한해 두었습니다. 결국, input과 output으로 가능한 데이터 타입들에는 다음이 있습니다:

int, float, string / PackedFunc / Module / DLTensor / TVM Object(IR의 모든 object들)

tvm.runtime.rpc

PackedFunc를 위한 RPC(Remote Procedure Call)입니다. cross-compile을 하거나, autotvm 등에서 퍼포먼스 기록을 위해 사용됩니다.

⚙️ tvm/ir

IRModule / Type / PassContext, Pass / Op 등으로 구성되어 있고, IR(Intermediate Representation) 함수(ex. relay.Function, tir.PrimFunc)를 위한 자료구조 및 인터페이스를 가지고 있습니다.

비유를 하자면, 나무와 숲이 있을 때, 숲을 relay.Function, 나무를 tir.PrimFunc로 비유할 수 있습니다.

relay.Function

high-level functional program representation으로, end-to-end model에 해당한다고 볼 수 있습니다. control-flow, recursion, 이외 다양한 자료구조들이 포함된 computational graph라고 보면 됩니다.

tir.PrimFunc

low-level functional program representation으로, loop-nest choice, multi-dimensional load/store, threading, vector/tensor instruction 등을 포함합니다. 모델에서 각 layer를 어떻게 실행시키는지를 알려줍니다.

relay.Function, tir.PrimFunc 등이 담고 있는 내용은 다르지만, 둘 다 IRModule을 바탕으로 하기 때문에 calling convention이 명확합니다. 덕분에 호출 등이 용이하고, 추후에는 cross-function-variant-optimization까지 가능해집니다.

PassContext는 pass 파이프라인을 가능하게 합니다. pass는 컴파일러의 pass를 의미합니다. 구체적인 내용은 Pass Infrastructure에서 확인 가능합니다.

Op는 모든 system-defined primitive operator/intrinsic을 나타내기 위한 클래스입니다.

⚙️ tvm/relay

전체 모델을 나타내기 위한 high-level functional IR입니다. relay.transform에 다양한 최적화 방식이 정의되어 있습니다. 다양한 dialect가 있고, 이들은 각각 서로 다른 최적화를 수행합니다.(ex. QNN, VM, Memory etc...)

아래 링크들은 2.3. Relay 파트에 더 자세하게 적어놓았으니 참고하면 좋을 것 같습니다.

Introduction to Relay IR

Relay Operator Strategy

Convert Layout Pass

⚙️ tvm/tir

위에서 tir.PrimFunc에 대해 설명할 때 언급한 바와 같이, tir(tensor IR)은 low-level program representation을 가지고 있습니다.

⚙️ tvm/te

tensor expression의 약자입니다. te는 그 자체로 IRModule을 구성하기보다는, te 블록들의 조합으로 IRModule을 구성합니다.

DSL(Domain Specific Language) 모듈로, te를 이용하여 tir::PrimFunc을 빠르게 만들 수 있습니다.

⚙️ tvm/topi

Tensor OPerator Inventory입니다.

tir이나 te로 operator를 전부 나타내도 괜찮지만, 마치 어셈블리어로 코드를 작성하는 느낌입니다. topi에는 numpy로 작성된, 딥러닝에서 자주 사용되는 pre-defined operators(in te or tir)가 있습니다.

이를 바탕으로, 다양한 target platform에서 괜찮은 성능을 보이는 implementation을 위해 공통된 스케줄 템플릿을 제공하기도 합니다.

⚙️ relay? ir? tir? te? topi?가 헷갈린다면

Relay

graph optimization toolset입니다. te를 바탕으로 작동합니다.

Input te는 임포트된 trained neural network나 relay.<op>로 직접 작성한, Relay IR(=high-level IR)입니다.

각각의 relay operator들은 어트리뷰트를 갖고 있습니다. 이들은 나중에 FTVMStrategy를 사용하여 최적의 topi implementation을 찾는데 사용됩니다.

tir

target-specific 백엔드로 translate되기 전 모습의 low-level IR입니다.

예를 들어, target이 x86이라면, tir은 llvm IR로 매핑되고 나서 target binary를 위해 x86 codegen이 호출되는 형식입니다.

topi

Relay IR과 tir 사이의 가교 역할을 수행합니다. operator들이 있을 때, 무엇을/어떻게 연산해야 하는지 정의하는 레퍼런스라 생각해도 좋습니다. FTVMSTrategy 를 통해 relay operator들을 highest scored implementation으로 매핑하도록 해 줍니다.

relay와 tir의 차이점이라면, Relay는 graph optimization toolset, 즉 전체적인 full-program-level optimization을 수행하는 반면 TIR은 dense linear algebra kernel와 같이 loop optimization을 목적으로 한다는 점에 있습니다.

relay에서 tir로의 변환은 relay.backend.compile_engine.py의 CompileEngine 클래스에 있는 lower 함수에서 이루어집니다.

⚙️ tvm/autotvm

AutoTVM의 작동 원리에 대해서는 여기를 보시면 더 자세한 설명을 볼 수 있습니다.

AutoTVM은 AutoScheduler와 함께 search based program optimization을 자동화시킵니다. 크게 3가지 역할을 수행하는데,

cost model, feature extraction

record format: Autotvm Database

search policy: Autotvm Exploration Module
가 있습니다.

topi와 autotvm의 관계에 대해서는 다음의 글에 잘 나타나 있습니다.

2.3 Relay

필자의 프로그래밍언어론에 대한 배경지식이 부족하여, relay에 대한 설명은 논문에 기반한 원리 설명보다는, design and architecture 부분에 관련한 내용들로 채웠습니다. 그래도 TVM 코드를 이해하는 데에는 많은 도움이 되는 글들입니다!

TVM 초반에는 NNVM 프로젝트로 불렸다가, 현재는 Relay로 완전히 대체되었습니다.

아래는 Design-and-Architecture 페이지에 나오는 relay에 대한 추가설명을 정리한 글입니다.

📑 Introduction to Relay IR

📑 Relay Operator Strategy

📑 Convert Layout Pass

아래 tutorial 파트에서 막힌 부분이 있는데, relay 부분을 다시 찬찬히 살펴본다면 해결할 수 있을 것 같은 문제로 보입니다.

🔎 함께 곁들이면 좋은 문헌:

2.4. AutoTVM

AutoTVM이란

AutoTVM은 머신러닝을 사용하여 딥러닝의 tensor operation들을 자동으로 최적화시켜주는 프레임워크입니다.

우리의 코드가 원하는 하드웨어에서 잘 돌아가기 위해서는 tensor operation library가 필요합니다. 예를 들면 GPU의 cuDNN, 인텔 프로세서의 MKL 등이 있습니다. 하지만, 이러한 tensor operation library들은 지원하는 하드웨어의 범위가 넓지 않고, 또 manually optimized 되었다는 단점을 가지고 있습니다.

💡 Manually optimized?

하드웨어마다 그 특성이 다르기에 threading/memory/pipelining 등등 다양한 변수들이 존재합니다. 이때, tensor operation library를 만드려면 이들 요소를 전부 고려해야 하는데, 이를 엔지니어들이 전부 휴리스틱하게 작성했다는 의미입니다.

이는 너무나도 많은 비용이 드는 방식이고, 다양한 딥러닝 프레임워크와 하드웨어들이 새롭게 등장하고 발전하는 현 시점에 각각의 조합에 맞는 방법들을 일일이 만드는 것은 비효율적입니다.

AutoTVM의 원리

AutoTVM은 "statistical cost model"을 만들어 주어진 코드의 runtime을 예측할 수 있도록 하고, 예측한 runtime을 바탕으로 최적화를 자동화합니다.

ninety

출처:earning to Optimize Tensor Programs

: index expression (ex. )
: transformation(index expression -> low level code)
: 주어진 에 대한 가능한 모든 의 집합
: 컴파일러(low-level code 생성기)
: 컴파일러를 통해 index expression에서 low level code로 변환
: 생성된 low-level code
: low-level code의 실행시간

위 표현을 사용해 우리의 목적을 나타내어 보자면, 아래의 식을 만족하는 를 찾는 것입니다:

AutoTVM의 전체적인 파이프라인은 다음과 같습니다.

statistical cost model 는 low level program의 cost를 예측하는 부분이고, exploration module에서는 다음에 어떤 schedule 를 적용할지 결정합니다.

이 schedule을 직접 하드웨어에서 실행시키고, 실행 결과들(시간이 얼마나 걸렸는지 등)은 데이터베이스 에 저장됩니다. 이들은 다시 Cost Model의 예측 정확도를 강화시키는 데 사용됩니다.

AutoTVM의 구성 요소

Statistical Cost Model
- (low-level abstract syntax tree(AST))가 주어졌을 때, 이의 domain-specific feature를 추출합니다. domain-specific feature라 함은 loop structure information(memory access 횟수, 데이터 재사용 비율 등)과 기타 정보들(vectorization, unrolling, thread binding 등)을 포함합니다.
  
  objective function은 아래 수식인 rank loss function으로 잡고 Database 의 결과들을 바탕으로 cost model을 학습시킵니다.
  
  Rank loss function:
- 논문에서는 XGBoost와 TreeGRU를 사용한다고 합니다.
  
  (하지만 tvm 오픈소스에서는 TreeGRU를 사용하는 부분이 보이지 않는군요. TVM discuss 홈페이지에서 누군가가 질문했고, 아직 merge되진 않았지만 관련 코드는 있습니다.)
Exploration Module
페이퍼에 적힌 알고리즘을 보면서 이해하는 것이 좋을 것 같습니다.

full

출처 : [Learning to Optimize Tensor Programs(https://arxiv.org/abs/1805.08166)

알고리즘을 보면 가능한 모든 를 살펴보기에는 search space가 너무 크기 때문에, 가장 좋은 효율을 보여줄 곳이라고 기대되는 부분들을 집중적으로 탐색합니다.(상단의 Q와 S를 다루는 부분) 우선 를 energy function으로 잡고 simulated annealing를 바탕으로 candidate Q를 설정합니다. -greedy 알고리즘을 사용하여 exploration을 보장하는데, 만큼은 랜덤하게 샘플하고, 나머지 만큼은 greedy하게 Q에서 다음의 목적함수를 maximize시키도록 합니다:

앞의 항은 runtime이 적은 를 고르기 위함이고, 마지막 항은 가 얼마나 많은 configuration component를 커버할 수 있는가를 나타냅니다.

이외에도 여러개의 workload에 대해 빠르게 해답을 찾을 수 있도록 transfer learning을 활용하는 등, 다양한 내용들과 이에 대한 실험결과들이 나와있는데, Learning to Optimize Tensor Programs를 참고하면 좋을 것 같습니다!

🔎 함께 곁들이면 좋은 문헌:

Learning to Optimize Tensor Programs

2.5. VTA

VTA란

도큐먼트를 보면 VTA를 다음과 같이 정의합니다:

💡 VTA (versatile tensor accelerator) is an open-source deep learning accelerator complemented with an end-to-end TVM-based compiler stack.

의역하자면, VTA는 TVM의 end-to-end compiler에서 하나의 end를 맡는, 오픈소스 deep learning accelerator 스택입니다.

eighty

출처: https://tvm.apache.org/docs/vta/dev/index.html

VTA의 특징

Generic, modular, open-source hardware
FPGA에서 사용하기 용이하고, FPGA가 없더라도 사용할 수 있도록 simulator를 제공합니다.
Simulator/FPGA를 위한 pynq 기반 드라이버 및 JIT(Just In Time) runtime을 제공합니다.
- bitstream.py 파일을 보면 FPGA에 bitstream을 다운받고 굽습니다. vta.program_fpga()한 줄이면 충분합니다!
TVM과 함께 쓰기 용이합니다.
- TVM을 통해 deep learning framework의 모델을 직접적으로 optimize하고 deploy할 수 있습니다.
- RPC(Remote Procedure Call)을 통해 deploy하기 쉽고, 또 파이썬을 통해 FPGA를 프로그래밍할 수 있습니다.

VTA Architecture

eighty

출처: https://tvm.apache.org/docs/vta/index.html

MODULES

크게 4가지 모듈[fetch, load, compute, store]로 구성되어 있습니다. 이들은 서로 SRAM buffer를 통해 TLP(task-level parallelism)을 가능하게 합니다.

fetch 모듈은 DRAM에서 instruction stream을 가져와서, 그에 맞게 command queue에 명령을 보냅니다.
load 모듈은 input/weight 데이터를 DRAM에서 가져와 on-chip memory로 배치시켜줍니다.
compute 모듈은 GEMM core와 ALU로 이루어져 있는데, GEMM core는 dense linear algebra computation을 수행하고, ALU는 나머지 기타 연산을 수행합니다. 연산 뿐 아니라, load 모듈처럼 적절한 장소에 데이터를 배치하는 업무도 수행합니다.(DRAM에서 레지스터로 데이터 가져오기, micro-op kernel을 micro-op cache로 가져오기 등등)
store 모듈은 연산이 완료된 결과를 DRAM에 다시 보내주는 역할을 수행합니다.

이들 모듈에 관한 자세한 설명은 여기에서확인할 수 있습니다.

VTA ISA

VTA는 4가지로 구성된, 독자적인 CISC 형태의 ISA(Instruction Set Architecture)를 가지고 있습니다.

LOAD Instruction: 데이터와 명령어를 준비하는 역할을 수행합니다(Load 모듈에서 수행하는 buffer/register로의 데이터 배치와, micro-op cache에 micro-kernel 넣기 등등)
GEMM Instruction: 행렬-행렬 곱셈을 수행하고 레지스터에 계산된 텐서값을 넣습니다.
ALU Instruction: 레지스터에 있는 텐서에 대해 ALU 연산을 수행합니다.
Store Instruction: output buffer에 있는 2D 텐서를 DRAM에 저장합니다.
ISA의 구조는 아래와 같고, 추가적인 정보는 다음의 문서에 나와있습니다.

ninety

출처: https://tvm.apache.org/docs/vta/dev/hardware.html#microarchitectural-overview

Details

각각의 모듈 간 존재하는 FIFO queue/buffer를 통해 모듈들이 서로 동기화할 수 있고, concurrent한 작업들을 수행할 수 있습니다. 덕분에 dataflow fashion으로 작업을 수행할 수 있습니다.

ninety

출처: https://tvm.apache.org/docs/vta/dev/hardware.html#microarchitectural-overview

이러한 dataflow execution을 바탕으로 현재는 3-stage pipelining(load-compute-store)이 구현되어 있습니다. superpipeline처럼 더 많은 stage를 가질 수 있도록 확장시키는 것은 logic overhead가 발생할 수 있기에 VTA는 3-stage pipeline으로 작업을 수행한다고 합니다.

개인적인 생각으로, MLPerf에서 VTA를 한번 테스트해보면 재미있는 프로젝트일 것 같네요(아직은 개발중인 프로젝트라 MLPerf를 수행하기는 힘든 것 같습니다).

🔎 함께 곁들이면 좋은 문헌:

3. Tutorial: Implementing Cartpole at Pynq-Z1 board with TVM

TVM tutorial 페이지에서 TVM을 활용한 다양한 튜토리얼 예시들을 볼 수 있습니다. 혹은, 다음의 페이지에서도 다양한 TVM tutorial을 확인할 수 있습니다.

이번 포스트에서는, 전술한 페이지에 없는 새로운 task를 진행해 보고자 합니다. DQN을 활용해 Cartpole 문제를 Pynq-Z1 보드에서 해결하는 것이 문제입니다.

(⚠️ AutoTVM, VTA를 전부 이용해 완벽하게 구현하려 하였으나, 시간 관계상 다 마치지 못하였습니다. 이에 ~~난중일기처럼~~ 회상록 형식으로 작성하겠습니다.)

Cartpole의 코드는 미디엄의 글에 올라온 코드를 살짝 수정하였습니다. 초반에는 파이토치 튜토리얼에 있는 코드를 사용하다가, 이 코드가 더 간단하고 수정하기 용이하여 사용하였습니다.

3.1. Environment

Xilinx Pynq-Z1 board with v2.5 image, 64GB SD Card
Python version 3.6.9
Ubuntu 18.04 with VirtualBox at MacBookPro19 with macOS Big Sur 11.1
- Network setting:
  
  Adapter 1: Intel PRO/1000 MT Desktop (NAT) - Disable
  
  Adapter 2: Intel PRO/1000 MT Desktop (Bridged Adapter, en6: Ethernet Adapter) - Enable
Development Platform: Visual studio code
Pynq Connection Environment
- Pynq에서 인터넷을 사용하고, 호스트 PC와의 IP 주소 동기화를 위해, iptime 공유기를 사용하였습니다.
VirtuaBox를 사용하는 경우, Network 설정을 잘 해 주어야 합니다. 기본값인 NAT로 설정하는 경우 VTA&AutoTVM을 사용할 때 문제가 생길 수 있습니다. 또한, LAN케이블로 연결한 것과 동일한 Wifi로 접속해야 VirtuaBox Ubuntu로의 ssh 접속이 가능합니다.
저의 경우, Bridged Adapter-이더넷 어댑터를 사용하였습니다. 각 네트워크 구성에 대한 보다 자세한 내용은, 다음 페이지를 참고하였습니다.
접속할 때, xilinx@192.168.2.99로 접속하는 대신 와이파이 공유기에서 할당한 내부 ip로 네트워크를 연결하였습니다. 공유기 주소 192.168.0.1로 접속하였을 때 네트워크는 다음과 같이 구성되어 있음을 알 수 있습니다.

ninety

VirtualBox에서 pynq로의 접속은 xilinx@192.168.0.36을 통해 진행하였습니다. 접속이 잘 되는지의 여부는, VirtualBox에서 ssh로 pynq에 접속해 보거나, 반대로 pynq에서 VirtualBox로 접속해 보는 방법이 있습니다. 상호 ssh 접속을 위해 authenticity도 설정해 주어야 하니, 직접 ssh로 접속하여 확인해 보는 것을 추천드립니다.
공유기 내부 IP는 pynq와 host 머신에 대해서는 수동으로 할당해 주는 것이 좋습니다

ninety

3.2. Memoir & TroubleShootings

PYNQ-Torch

Pynq에서 강화학습 문제를 해결하는 것이 첫번째 목표였습니다. 처음에는 PYNQ-Torch를 사용하였습니다. PYNQ의 하드웨어적 사양은 몹시 낮습니다. CPU는 갤럭시S2와 동급인 것은 둘째치고, 특히 RAM은 kB 단위입니다. RAM 용량이 몹시 낮기 때문에 일반적인 pytorch를 Pynq 보드에서 작동시키기는 힘듭니다. 이에, Pynq의 이미지가 들어있는 SD 카드의 일부분을 Swap-Memory 구역으로 지정하여 RAM처럼 사용하도록 하게 하는 것이, Pynq-Torch입니다.

Pynq-Torch를 사용하기 위해서는 SD 카드의 용량이 64GB는 되어야 합니다. Pynq-Torch 설치 및 다양한 튜토리얼들은 PYNQ-Torch 깃허브에 있습니다. Pynq의 CPU는 몹시 느리기 때문에, 설치도 느립니다.
- https://github.com/manoharvhr/PYNQ-Torch
- https://www.researchgate.net/publication/337030249_PYNQ-Torch_a_framework_to_develop_PyTorch_accelerators_on_the_PYNQ_platform
Cart-pole 문제는 잘 해결되었지만, 단순히 작동한다는 것이지, 주어진 Zynq SoC 사양을 최대한으로 활용할 수 있도록 최적화를 잘 해주는가에 대해서는 회의적이었기에, Automated, End-to-End Compiler인 TVM을 적용하여 Cartpole 문제를 해결하기로 마음먹었습니다.
TVM 설치 - Troubleshooting

생각보다 TVM을 설치하는 과정이 힘들었습니다.
1. llvm path 문제
  
  📑 Cross Compilation and RPC - tvm 0.8.dev0 documentation
  
  위 포스트처럼 pynq에서는 최대한 가벼운 tvm runtime만 설치하고, 나머지 optimization의 역할을 수행하는 host 부분은 Mac에서 작동시키려 하였습니다.
  
  원래는 Mac OsX에서 tvm을 설치하려 하였으나, llvm의 경로를 잘 찾지 못해 빌드하는 과정에서 계속 실패하였습니다.
  
  결국 virtualbox에 ubuntu를 설치하고, apt-get를 사용하여 llvm을 설치하니 경로 문제가 단번에 해결되었습니다. 빌드에 있어서는 리눅스 OS가 편한 것 같습니다.
  
  sudo bash -c "$(wget -O - https://apt.llvm.org/llvm.sh)"
  
  로 llvm을 설치한 후,
  
  llvm-config --version
  
  커맨드로 llvm 버전을 확인하면 됩니다.
  
  llvm 설치 후 tvm 설치과정 중 config.cmake 파일에서 llvm 설정을
  
  set(USE_LLVM ON)과 set(LLVM-CONFIG llvm-config-본인 버전)로 입력하면 됩니다. 나머지 부분은 본인의 입맛에 맞게 ON과 OFF를 잘 설정해 주세요.
2. VTA RPC 과정에서 Recursion Error가 발생하였습니다.
  
  https://discuss.tvm.apache.org/t/vta-recursion-error/7271
  
  빌드의 문제였고, make clean 후 다시 빌드하니 잘 되었습니다.
3. Install from source가 가장 쉬운 방법인 듯 합니다. MacOS에서 Dockerfile로 install을 시도하였을 때 포트 관련한 에러가 났습니다.
Implementing Cartpole with TVM

아쉽지만 시간 관계상 미완성의 상태로 마무리하게 되었습니다. 본래 목적은 ARM CPU+FPGA의 Zynq SoC를 가진 pynq 보드에 DQN을 이용하여 cartpole을 완전히 해결하는 것이었습니다.

현재는 ARM CPU에서만 작동하는 코드까지 완성되었습니다. VTA와 AutoTVM을 활용하여 FPGA까지 활용하는 것이 최종 목표입니다. 지금까지의 코드를 다음 깃허브에 올려두었습니다.
- ARM CPU에서 Cartpole 실행하기 (main branch)
  - https://github.com/kc-ml2/tvm
- VTA, AutoTVM을 사용, CPU+FPGA를 완벽하게 활용하기 (vta_cartpole branch)
  - https://github.com/kc-ml2/tvm/tree/vta_cartpole

4. Discussion and Summary

TVM 번역 작업이 진행중입니다. 관심 있으신 분은 함께해 주시면 감사하겠습니다.

Reference

TVM Homepage
TVM Source Codes
TVM Conference
TVM: An Automated End-to-End Optimizing Compiler for Deep Learning
TVM: End-to-End Optimization Stack for Deep Learning
Relay: A High-Level Compiler for Deep Learning
Learning to Optimize Tensor Programs
A Hardware-Software Blueprint for Flexible Deep Learning Specialization
Pytorch acceleration in Pynq
CK-tvm
- https://github.com/ctuning/ck-tvm

이우진

Woojin LEE

Internship

0. Disclaimer

1. Introduction to TVM

2. TVM Workflow and Components

출처: https://tvm.apache.org/

2.1. Overall Workflow

Example Compilation Flow.(출처 : https://tvm.apache.org/docs/dev/index.html)

Transformation

Translation

2.2. Logical Components

TVM Architecture Design (출처: <"https://tvm.apache.org/docs/dev/index.html">)

⚙️ tvm/runtime

TVM supports multiple programming languages for the compiler stack development and deployment.

⚙️ tvm/ir

⚙️ tvm/relay

⚙️ tvm/tir

⚙️ tvm/te

⚙️ tvm/topi

⚙️ relay? ir? tir? te? topi?가 헷갈린다면

⚙️ tvm/autotvm

2.3 Relay

2.4. AutoTVM

AutoTVM이란

AutoTVM의 원리

출처:earning to Optimize Tensor Programs

AutoTVM의 구성 요소

출처 : [Learning to Optimize Tensor Programs(https://arxiv.org/abs/1805.08166)

2.5. VTA

VTA란

출처: https://tvm.apache.org/docs/vta/dev/index.html

VTA의 특징

VTA Architecture

출처: https://tvm.apache.org/docs/vta/index.html

출처: https://tvm.apache.org/docs/vta/dev/hardware.html#microarchitectural-overview

출처: https://tvm.apache.org/docs/vta/dev/hardware.html#microarchitectural-overview

3. Tutorial: Implementing Cartpole at Pynq-Z1 board with TVM

3.1. Environment

3.2. Memoir & TroubleShootings

4. Discussion and Summary

Reference