[kernel] Recompilation optimization triggered by triton function para…

[cvSoldier]

…meter optimization

What this PR does / why we need it?

Does this PR introduce any user-facing change?

no

How was this patch tested?

• vLLM version: v0.17.0
• vLLM main: vllm-project/vllm@8b63257

[gemini-code-assist]

Summary of Changes

Hello, I'm Gemini Code Assist¹! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!

This pull request implements an optimization strategy for Triton kernels within the vllm_ascend project. By removing the tl.constexpr annotation from various parameters in several kernel functions, the aim is to reduce the frequency of kernel recompilations. This adjustment allows the affected parameters to be passed dynamically, which can lead to performance improvements by avoiding redundant compilation steps for different parameter values, thereby streamlining the execution of these operations.

Highlights

• Triton Kernel Parameter Optimization: Modified several Triton kernel function signatures by removing the tl.constexpr qualifier from specific parameters.
• Improved Recompilation Efficiency: This change is intended to optimize kernel recompilation, likely by allowing these parameters to be dynamic without triggering a new compilation for each unique value.

🧠 New Feature in Public Preview: You can now enable Memory to help Gemini Code Assist learn from your team's feedback. This makes future code reviews more consistent and personalized to your project's style. Click here to enable Memory in your admin console.

Using Gemini Code Assist

The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.

Invoking Gemini

You can request assistance from Gemini at any point by creating a comment using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands on the current page.

Feature	Command	Description
Code Review	/gemini review	Performs a code review for the current pull request in its current state.
Pull Request Summary	/gemini summary	Provides a summary of the current pull request in its current state.
Comment	@gemini-code-assist	Responds in comments when explicitly tagged, both in pull request comments and review comments.
Help	/gemini help	Displays a list of available commands.

Customization

To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.

Limitations & Feedback

Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for GitHub and other Google products, sign up here.

Footnotes

1. Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩

[gemini-code-assist]

Code Review

This pull request optimizes Triton kernels by changing several tl.constexpr parameters to runtime arguments, which avoids kernel recompilation when these parameters change. The changes are applied across solve_tril.py, split_qkv_rmsnorm_mrope.py, and causal_conv1d.py. While the changes are generally correct and follow the intended optimization, I've found a critical issue in causal_conv1d.py where not all dependent parameters were updated, which will likely lead to a compilation error.

[gemini-code-assist]

While you've correctly updated many parameters to be runtime values to avoid recompilation, stride_w_dim and stride_w_width were missed. In the causal_conv1d_update_npu wrapper, stride_w_width is set to dim, which is now a runtime argument. Passing a runtime value to a tl.constexpr parameter will cause a compilation error. Both stride_w_dim and stride_w_width should be made runtime arguments for correctness and consistency.

Suggested change

	stride_x_seq,
	stride_x_dim,
	stride_x_token,
	stride_w_dim: tl.constexpr,
	stride_w_width: tl.constexpr,
	stride_x_seq,
	stride_x_dim,
	stride_x_token,
	stride_w_dim,
	stride_w_width,

[github-actions]

👋 Hi! Thank you for contributing to the vLLM Ascend project. The following points will speed up your PR merge:‌‌

• A PR should do only one thing, smaller PRs enable faster reviews.
• Every PR should include unit tests and end-to-end tests ‌to ensure it works and is not broken by other future PRs.
• Write the commit message by fulfilling the PR description to help reviewer and future developers understand.

If CI fails, you can run linting and testing checks locally according Contributing and Testing.