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

• [x] Python API workspace parameters {Turns, GPU number}: which is used by decorators to store temporary variable. [siahuat]
• [x] Kernel turning in sync mode [jilong]
  • [x] Invoke Antares using CLI
• [x] Kernel turning process display [weicui]
  • [x] python code to filter the process from Antares log
• [x] NNFusion Python wheel package [wenxiang]
• [x] Install Antares with NNFusion Python package
  • [x] Provide prebuilt package for specific env: coda, ... etc
• [ ] [Showcase] Fully automation of integration of existed workload: Install -> Running
  • [x] 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)

• [x] Memory manage feature {Pytorch allocation}
  • [x] nnfusion return a total workspace size: use a function in codegen result
  • [x] Change memory allocator in codegen source code
  • [x] 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
  • [x] Antares fusion @jlxue | 2days

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

  • [ ] Revisit operator set: Fix Conv, dot operator definition @xysmlx + other
  • [x] Script to run ORT optimization; @xysmlx #300
  • [ ] Script to run tf2onnx to convert tf pb to onnx format; @mzmssg
  • [x] 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 #50
  • [ ] Process __shared and syncthreads() with profiler to support any BlockCudaEmitter::get_kernel_cache_entry() function