Skip to content

Optimize pow function the focal loss,#1979

Closed
JohnNikolay84 wants to merge 310 commits intoNVIDIA:masterfrom
ROCm:focal_loss_pow_optim_prod
Closed

Optimize pow function the focal loss,#1979
JohnNikolay84 wants to merge 310 commits intoNVIDIA:masterfrom
ROCm:focal_loss_pow_optim_prod

Conversation

@JohnNikolay84
Copy link
Copy Markdown

@JohnNikolay84 JohnNikolay84 commented Jan 21, 2026

Pow implementation is very expensive on AMD mi355x. This commit changes it to a mathematically equivalent
exp(y*log(x)) for x > 0. However 1-2 ULP prec loss might be possible. Improvement however is very solid, the cost of the runtime for focal_loss_forward_cuda_kernel is more than 50% down.

image

lcskrishna and others added 30 commits August 18, 2020 11:53
* 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
* fix warp size in WARP_SHFL* in layernorm

* enable fused_layer_norm tests on ROCm
Hipify revamp changes for apex extensions on ROCm.
Fix reduce_block_into_lanes for multi_tensor_l2norm for ROCm
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
Mostly whitespace or formatting issues addressed.
Diff with upstream is reduced; ROCm changes are more clear.
use __launch_bounds__(1024) for multi_tensor_apply, re-enable skipped tests
- incorrect use of __shfl_down
- fix warp size assumptions
- update unit tests to exit on failure
Revert "pass all TensorListMetadata as pointer to pinned host memory (#13)
amd-sriram and others added 28 commits April 26, 2025 10:44
…port different grad (#207)

* Fix `DistributedFusedAdam` for grad dtype != param dtype (NVIDIA#1893)

* Pipeline `reduce-scatter` and `all-reduce`. (NVIDIA#1895)

---------

Co-authored-by: Tailing Yuan <yuantailing@gmail.com>
Co-authored-by: Wil Kong <alpha0422@gmail.com>
The error:
File /tmp/easy_install-_pfhn8pn/matplotlib-3.5.1/.eggs/setuptools_scm-8.3.1-py3.12.egg/setuptools_scm/_integration/pyproject_reading.py, line 36, in read_pyproject section = defn.get(tool, {})[tool_name] ~~~~~~~~~~~~~~~~~~~~^^^^^^^^^^^ KeyError: 'setuptools_scm'

Solution : https://github.com/matplotlib/matplotlib/blob/v3.8.x/pyproject.toml#L22 matplotlib 3.8 is the first version to have pyproject.toml with this tool.setuptools_scm section. This higher version of setuptools expects this structure in the python packages it installs. Matplotlib 3.5.1 doesn't satisfy this condition. The solution is to change the condition to matplotlib>=3.8.
…reductions code inside cuWelfordMuSigma2() function in layer norm kernel assumes a warp size of 32, so added a condition for rocm to support gpu warp size (based on earlier apex code). For rocm, adjust the threadsize, based on earlier apex code. (#215)
* Fix fused_dense_gelu_dense, change the names of the parameters so that they can be accessed by the test appropriately

* Update the absolute tolerances in test_mlp from 0 and 1e-7 to 1e-5

* Deactivate the amp state handle for optimization level other than O0. This helps to pass the UT after this.

* Update condition for deactivating amp state handle from opt level equal to 1 to opt level not equal to 0

* Update torch set default dtype method to remove warning

* Update the method to create overflow buffer for amp optimizer

* Update the method to create overflow buffer for amp optimizer

* Update the method to create overflow buffer for amp optimizer

* reset the default device to cpu so that the generator uses cuda, as run_amp tests set its to cuda
In ROCm 7.0, the warpSize variable is no longer constexpr.
This commit replaces the variable use with the correct values
based on the architecture we're running on.
Do not use warpSize as a constexpr in nhwc_batch_norm_kernel.h
* Added aiter support in fused_rope.py for all 4 variants. Updated fused rope test, reduced tolerances according to unit test in aiter repo.

* Add aiter as a submodule and install it if it is rocm. Switch on aiter backend if it is rocm and aiter is installed

* add pandas to the requirements so that aiter can be used without numpy error - ValueError: numpy.dtype size changed, may indicate binary incompatibility. Expected 96 from C header, got 88 from PyObject

* Replace ROCM_HOME condition to IS_ROCM_PYTORCH for installing aiter and use pip install -e . instead of python setup.py develop for installing aiter.

* Create apex and aiter subclasses for the four variants of FusedRoPEFunc and select apex or aiter subclass based on AITER_ROPE_BACKEND value. The user can specify the environment variable USE_ROCM_AITER_ROPE_BACKEND to select between aiter and apex backends for fused rope.

* If the AITER backend is selected, use lowered precision in the unit test otherwise use the original precision 1e-3

* warn user about the lower precision when using aiter backend for fused rope

* Update fused_rope.py

remove spaces

* simplify the switch between aiter and apex subclasses

* install aiter without editable mode
* add test to extract extensions from setup.py and test if there can be imported

* moved test outside tests/L0
* made a flag to switch on/off aiter compile using --aiter when installing apex

* Added information on building AITER during installation in readme
* replace c10_warp_size in fused rope

* replace c10_warp_size in fused softmax

* replace c10_warp_size in group batch norm

* replace c10_warp_size in multiheadattention

* replace c10_warp_size in tramsducer

* replace c10_warp_size in xentropy

* replace c10_warp_size in sync batch normalization

* replace c10_warp_size in group batch norm

* replace warp_size in multihead attention
This commit changes it to a mathematically equivalent
exp(y*log(x)) for x > 0.
However 1-2 ULP prec loss might be possible.
@amd-sriram amd-sriram deleted the focal_loss_pow_optim_prod branch January 22, 2026 09:41
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.