[feat] Add causal_conv1d_update triton kernel

[OsirisDuan]

What this PR does / why we need it?

This PR introduces a new Triton kernel _causal_conv1d_update_kernel_no_cache_len_no_mtp to support efficient causal 1D convolution updates in Qwen3-next. The kernel is integrated into causal_conv1d_update_npu, enabling better performance on Ascend NPU hardware.

Does this PR introduce any user-facing change?

It belongs to optimization at the internal implementation level, which helps improve the final model inference performance but does not change the user experience at the API level or in terms of invocation methods.

How was this patch tested?

• vLLM version: v0.11.2
• vLLM main: https://github.com/vllm-project/vllm/commit/v0.11.2

[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.

[gemini-code-assist]

Code Review

This pull request introduces a new Triton kernel _causal_conv1d_update_kernel_no_cache_len_no_mtp for a specific case of causal 1D convolution, and integrates it into the existing causal_conv1d_update_npu function. While the addition is a good optimization, I've found a critical bug in the implementation of the new kernel where it fails to correctly write outputs for sequences longer than one token, causing data loss. My review includes a specific comment with a code suggestion to fix this issue.

[gemini-code-assist]

There is a critical bug in this loop. The tl.store operation writes to a memory location that does not depend on the loop variable i. As a result, for sequences with seq_len > 1, the output for each token will be written to the same location, overwriting the previous one. Only the result for the last token (i = seq_len - 1) will be preserved, leading to incorrect convolution output. To fix this, you need to include the token index i in the output pointer calculation to ensure each token's result is stored in its correct position.

Suggested change

	tl.store(
	out_ptr
	+ pid * out_batch_stride
	+ (doffs + tl.arange(0, DIM_BLOCK)) * out_len,
	result,
	)
	tl.store(
	out_ptr
	+ pid * out_batch_stride
	+ (doffs + tl.arange(0, DIM_BLOCK)) * out_len + i,
	result,
	)

[github-actions]

This pull request has conflicts, please resolve those before we can evaluate the pull request.