NNFusion v0.4 (JUN~) Release Plan

[wenxcs]

Release Plan for V0.4

• Release Manager: TBD
• Feature Freeze Date: TBD
• Code Freeze Date:
• Test Verification/Bug Bash:
• Release Date:
• Release Note:

---

Annotation Legend:

Mark	Description
🏃‍♂️	work in progress
👪	work under review
✋	blocked task

Feature

Python API

• Python API workspace parameters {Turns, GPU number}: which is used by decorators to store temporary variable. [siahuat]
• Kernel turning in sync mode [jilong]
  • Invoke Antares using CLI
• Kernel turning process display [weicui]
  • python code to filter the process from Antares log
• NNFusion Python wheel package [wenxiang]
• Install Antares with NNFusion Python package
  • Provide prebuilt package for specific env: coda, ... etc
• [Showcase] Fully automation of integration of existed workload: Install -> Running
  • import fail for resnet50.fp16 model [Constant op]
  • import fail for densenet169.fp16 model[Pad op]
  • tensorcore FP16
  • Document

Customized operator [ziming]

• Shape inference (json field to put input/output shape) of customized operator in Python API
• Customized kernel code {not using Antares} in Python API

Function codegen

(Output tensor to use easiest way to allocate)

• Memory manage feature {Pytorch allocation}
  • nnfusion return a total workspace size: use a function in codegen result
  • Change memory allocator in codegen source code
  • Use PyTorch to allocate memory pool [siahuat]

Other

• Constant folding(use onnx to do this)
• Expose NNFusion log level option

[Not included] Code generator

In this section CG will represent our new code generator

• Integration of CG
• Tear down and reassemble CG as our purpose

Tensor Program level API

• Axis fusion
• Tensor Padding
• Expose Variable as input shape

Schedule

• Sparse schedule primitive: API to specify an axis as sparse loop, with a CSR data as input
• Sparse schedule primitive: Specify dense body and its data offset
• Primitive to unroll a sparse loop

Stretch

Mechanism

• Graph Rewriting Feature @wenxcs ? weeks |

  • Sub-graph match mechanism @xiayuqing0622
  • Operator replacing: replaced with non-existed op; @jlxue | 2days
  • Kernel replacing @wenxcs
  • Conv + element wise fusion @xiayuqing0622
  • Antares fusion @jlxue | 2days

• Enhance user experience and support for ONNX format @jlxue | 1week

  • Revisit operator set: Fix Conv, dot operator definition @xysmlx + other
  • Script to run ORT optimization; @xysmlx Automatically optimize ONNX Graph with external tool #300
  • Script to run tf2onnx to convert tf pb to onnx format; @mzmssg
  • Script to compare with ORT; @xysmlx

• Unify kernel provider interface for Antares, KernelDB, and our internal kernels @wenxcs 1 weeks |

  • Design general kernel provider interface;
    • Generate identifier to identify unique op configs | @xysmlx

Refactor/Improvement

• Element wise operator performance @xysmlx
• Constant folding performance @yiyione (TBD)
  -[ ] Profiler performance
  Kernel Performance:
  • RowReduceSum Kernel performance @xysmlx

• [Low priority] Serialization and deserialization interface
  • High level IR representation;

• docker image support low version of cuda driver @wenxcs 3 days |
• Investigation on GNN training @xysmlx
• Merge Gamefusion: reduce fusion @yiyione
• Benchmark Page @wenxcs 3 days |

• Quant&Sparity Interface
• CPU performance improvement
• Android code-gen&tuning
• Control-flow support ｜@xysmlx
  • python integration
• new profiler: support different profiling information, e.g., time, kernel time provided by nvprof | @xysmlx
  • Base profiler | 5 days
  • Update existing profilers | 1 day
  • CUDA profiler: kernel time, num of __syncthreads(), num of active thread blocks, shared memory, grid/block size | 2 days
  • Kernel profiling pass | 1-2 days
• Profiling cache DB | @xysmlx
  • DB schema and interfaces
  • Integrate with profiler
• BlockFusion | @xysmlx
  • BlockFusion active thread block check [ENHANCEMENT] Active block check in -fblockfusion_level=2 #50
  • Process __shared__ and __syncthreads() with profiler to support any BlockCudaEmitter::get_kernel_cache_entry() function