Skip to content

[MLU] feat(fused_kernel): add fused linear_cross_entropy kernel for MLU - #12

Open
uqyxx wants to merge 5 commits into
verl-project:mainfrom
uqyxx:mlu-fused-linear-cross-entropy
Open

[MLU] feat(fused_kernel): add fused linear_cross_entropy kernel for MLU#12
uqyxx wants to merge 5 commits into
verl-project:mainfrom
uqyxx:mlu-fused-linear-cross-entropy

Conversation

@uqyxx

@uqyxx uqyxx commented Aug 26, 2026

Copy link
Copy Markdown
Contributor

Summary

Add MLU fused linear_cross_entropy kernel.

Changes

Add linear_cross_entropy and function_wrapper to replace original linear_cross_entropy when running on MLU.

Testing

  • pytest tests/ -v passes
  • Manually verified on target hardware (if applicable)

Checklist

  • Code follows the project's style and passes pre-commit checks
  • Documentation updated (if applicable)
  • No secrets or credentials included

Comment thread verl_hardware_plugin/engines/fsdp_mlu.py Outdated
Comment thread verl_hardware_plugin/utils/kernels_mlu.py Outdated
Comment thread verl_hardware_plugin/utils/function_wrapper.py Outdated
Comment thread verl_hardware_plugin/utils/kernels_mlu.py Outdated
@uqyxx

uqyxx commented Sep 9, 2026

Copy link
Copy Markdown
Contributor Author

Hi @physics31415926 , thanks for your comments, all issues resolved, please review.

accu_ptrs = accu_ptr + offs_m_cur[:, None] * stride_accu_m + offs_n[None, :] * stride_accu_n
entropy_b_ptrs = entropy_b_ptr + offs_m_cur[:, None] * stride_entropy_b_m + offs_n[None, :] * stride_entropy_b_n

_max = tl.load(max_ptrs, mask=mask_n)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Please set other on these masked loads (-inf for maxima, 0 for accumulators). For vocab_size=151936, 149 splits occupy a 256-wide block, so undefined masked lanes enter the reductions and can corrupt logprobs/entropy. The TP epilogue needs the same fix.

pid_n = tile_idx // num_pid_m

if pid_m == 0 and pid_n == 0:
tl.store(global_logprobs_scalar_ptr, 0.0)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

With reduction="none", this pointer aliases logprobs[0]. If another program writes the first token's label logit before this store, it gets overwritten with zero. Please remove this store; sum/mean outputs are already zero-initialized on the host.

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.

2 participants