* [TOPI][CUDA] Improve the performance of scatter_nd by:
1. Split into 2 kernels, one does the "Init" and another does the "Update".
Thus they can have different Grid/Block configurations to better utilize
SMs.
2. Use atomic_add instead of direct assignment, which could avoid the race
condtion when multiple indices point to the same location of the output
tensor. With this moidification, it's safe now to use more CUDA threads
to gain more parallelism.
* Fix python code format.
* FIX: [TOPI][CUDA] Improve the performance of scatter_nd #8479
- Split ScatterND kernel into 2 sub-kernels using ib.new_scope()
- Replace ib.for_range() with blockIdx.y
- Using atomic_add when mode == "add"
- Keep threadIdx.x less than max_threads of GPU
* Comment added
* Add fallback implementation when "mode=add" meets int64
- Atomic_add from CUDA doesn't support int64 data type
- Change "ind{i}" to "ind%d"%i, where names of relay.var could correctly display
* Python format
* Fix line too long
* CI pass
* Empty, for CI pass
* Empty, for CI pass
* Empty, for CI pass
* Empty, for CI pass
* Empty, for CI pass
* Exchange blockIdx.x and blockIdx.y
* check for Vulkan or metal
* Fallback to previous algorithm when mode==update
* Update python/tvm/topi/cuda/scatter.py
Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com>
* Assign TODO
* Swapping then and else block
Co-authored-by: wenxizhu <wenxizhu@tencent.com>
Co-authored-by: CaptainDuke <captainduke328@gmail.com>
Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com>
zhuwenxi
887324f778