GitLab

use __launch_bounds__(1024) for multi_tensor_apply, re-enable skipped tests

IFU-2021-01-18

Mostly whitespace or formatting issues addressed.
Diff with upstream is reduced; ROCm changes are more clear.

Conflicts:
csrc/multi_tensor_apply.cuh
setup.py
tests/L0/run_optimizers/test_adagrad.py
tests/L0/run_optimizers/test_fused_optimizer.py
tests/L0/run_optimizers/test_lamb.py

Fix reduce_block_into_lanes for multi_tensor_l2norm for ROCm

Skip the unit tests

Update ASP README to highlight default recipe

The Recipe was presented after some non-standard API calls, so moving the suggested usage up, giving it its own section, and reinforcing the suggested usage in the non-standard section.

Hipify revamp changes for apex extensions on ROCm.

* add flag for DistributedAdam: step_support_amp_scaling
Co-authored-by: Kexin Yu <kexiny@nvidia.com>
Co-authored-by: Kexin Yu <kexinznzn›@gmail.com>

* fuse dropout into softmax in fprop for additive mask case

- resume() is a nested function and when it loads best_prec1
  it creates a local variable that hides the one from the parent
  function (which refers to the global one). This PR adds `global`
  to modify the global variable as intended
Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>

DistributedFusedAdam Model Parallelism Support (Megatron)
Co-authored-by: Kexin Yu <kexiny@nvidia.com>
Co-authored-by: Kexin Yu <kexinznzn›@gmail.com>

* fix warp size in WARP_SHFL* in layernorm

* enable fused_layer_norm tests on ROCm

In this PR, we mainly tried to optimize the performance of Syncatchnorm and also fixed one potential issue in the welford_parallel kernel implementation.

For performance improvement, we batched the mean/var/count all_gather communication together and sent it once in the forward path
We also batch the all_reduce in backward path
We add the contiguous call on the input of welford_parallel kernel.
If there is any standard perf benchmark, I would be happy to run it.

Update asp readme

Update README for ASP

Added an outline to illustrate our recommended recipe to obtain a pruned model

* enable deprecated fused adam optimizer

* enable deprecated fused lamb

* enable xentropy extension

* add warpsize 32 for nv and 64 for amd

* update compiler arguments

* update the syncwarp conditions

* update syncwarp condition

* enable deprecated fused adam optimizer

* enable deprecated fused lamb

* reset the compiler arguments

* syntax error

* aligning the compiler arguments

Co-authored-by: pbialecki <pbialecki@nvidia.com>

* enable mlp cuda

* add setup changes and tests

* skip the unit tests

* updated conditions for empty array

* removed hip platform conditions