Skip to content

Add custom multi_tensor_apply kernels (L2norm, Adam)#585

Draft
matthiasdiener wants to merge 2 commits into
devfrom
mdiener/multi_tensor_apply_kernel
Draft

Add custom multi_tensor_apply kernels (L2norm, Adam)#585
matthiasdiener wants to merge 2 commits into
devfrom
mdiener/multi_tensor_apply_kernel

Conversation

@matthiasdiener
Copy link
Copy Markdown
Contributor

Description

Fixes https://github.com/ROCm/frameworks-internal/issues/16529

Type of change

  • Documentation change (change only to the documentation, either a fix or a new content)
  • Bug fix (non-breaking change which fixes an issue)
  • New feature (non-breaking change which adds functionality)
  • Breaking change (fix or feature that would cause existing functionality to not work as expected)
  • Infra/Build change
  • Code refactoring

Changes

Please list the changes introduced in this PR:

  • Change A
  • Change B

Checklist:

  • I have read and followed the contributing guidelines
  • The functionality is complete
  • I have commented my code, particularly in hard-to-understand areas
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have added tests that prove my fix is effective or that my feature works
  • New and existing unit tests pass locally with my changes

@matthiasdiener matthiasdiener self-assigned this May 13, 2026
@matthiasdiener matthiasdiener added the ci-level 1 CI test level 1 label May 13, 2026
@matthiasdiener matthiasdiener changed the title Add a custom multi_tensor_l2norm_kernel Add a custom multi_tensor_apply kernels (L2norm, Adam) May 15, 2026
@matthiasdiener matthiasdiener changed the title Add a custom multi_tensor_apply kernels (L2norm, Adam) Add custom multi_tensor_apply kernels (L2norm, Adam) May 15, 2026
template <int N, typename T>
__device__ __forceinline__ void load_store_n(T *dst, const T *src,
int dst_offset, int src_offset) {
typedef typename std::aligned_storage<N * sizeof(T), N * alignof(T)>::type LT;
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We already have some load store functions that are optimal for rocm here -- Can we reuse them?

https://github.com/ROCm/TransformerEngine/blob/dev/transformer_engine/common/util/rocm_device_utils.cuh#L68-L115

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P.S., this also has some other utils that come in handy for us.


TRANSFORMER_ENGINE_TYPE_SWITCH_NON_FP8ONLY(
grad_dtype, grad_type,
if (mode == ADAM_MODE_0) {
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We can use TRANSFORMER_ENGINE_SWITCH_CONDITION here I think

#pragma unroll
for (int ii = 0; ii < CILP; ii++) {
if (MODE == ADAM_MODE_0) { // L2
r_g[ii] = r_g[ii] + (decay * r_p[ii]);
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

+= here for readability

param_dtype, p_type,
TRANSFORMER_ENGINE_TYPE_SWITCH_NON_FP8ONLY(
grad_dtype, g_type,
if (mode == ADAM_MODE_0) {
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same, can use TRANSFORMER_ENGINE_SWITCH_CONDITION here

LAUNCH_CUSTOM_ADAM(g_type, p_type, ADAM_MODE_0, true);
} else {
LAUNCH_CUSTOM_ADAM(g_type, p_type, ADAM_MODE_1, true);
};););
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Add NOLINT here and at the all of our macro switches

@@ -1 +1 @@
/*************************************************************************
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copyright

@@ -1,2 +1,2 @@
/*************************************************************************
* Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copyright

@@ -1,2 +1,2 @@
/*************************************************************************
* Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copyright

@@ -1 +1 @@
# Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copyright

multi_tensor_scale,
multi_tensor_l2norm,
multi_tensor_unscale_l2norm,
multi_tensor_scale as _multi_tensor_scale,
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These changes should be hip guarded I think

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ci-level 1 CI test level 1

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants