Skip to content

[BugFix]support fa3 qwen-vl rope #5869

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ckl117 merged 1 commit into from
Jan 5, 2026
+26 −20
Merged

[BugFix]support fa3 qwen-vl rope #5869

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 19 additions & 17 deletions custom_ops/gpu_ops/append_attn/gqa_rope_write_cache.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1377,23 +1377,25 @@ std::vector<paddle::Tensor> GQARopeWriteCacheKernel(

if (use_neox_rotary_style) {
if (rotary_dim == head_dim) {
gqa_rotary_qk_split_variable_qwen3<data_t>(qkv_out.data<data_t>(),
q.data<data_t>(),
k.data<data_t>(),
v.data<data_t>(),
qkv.data<data_t>(),
rotary_embs.data<float>(),
batch_id_per_token.data<int>(),
seq_lens_encoder.data<int>(),
seq_lens_decoder.data<int>(),
cu_seqlens_q.data<int>(),
cu_seqlens_k.data<int>(),
token_num,
num_heads,
kv_num_heads,
max_seq_len,
head_dim,
stream);
gqa_rotary_qk_split_variable_qwen3<data_t>(
qkv_out.data<data_t>(),
q.data<data_t>(),
k.data<data_t>(),
v.data<data_t>(),
qkv.data<data_t>(),
rotary_embs.data<float>(),
batch_id_per_token.data<int>(),
seq_lens_encoder.data<int>(),
seq_lens_decoder.data<int>(),
cu_seqlens_q.data<int>(),
cu_seqlens_k.data<int>(),
token_num,
num_heads,
kv_num_heads,
rope_3d ? rotary_embs.dims()[3] : rotary_embs.dims()[2],
head_dim,
rope_3d,
stream);
} else {
gqa_neox_partial_rotary_qk_split_variable<data_t>(
qkv_out.data<data_t>(),
Expand Down
10 changes: 7 additions & 3 deletions custom_ops/gpu_ops/append_attn/qwen3_rope.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,8 @@ __global__ void GQAVariableLengthRotarySplitKernel_Qwen3(
const int q_num_head,
const int kv_num_head,
const int max_model_len,
const int head_dim) {
const int head_dim,
const bool rope_3d) {
using LoadT = AlignedVector<T, VecSize>;
using LoadEmbT = AlignedVector<float, VecSize>;
LoadEmbT cos_emb_vec;
Expand Down Expand Up @@ -84,7 +85,8 @@ __global__ void GQAVariableLengthRotarySplitKernel_Qwen3(
}

// TODO check this correct or not
int64_t new_emb_idx = emb_idx;
int64_t new_emb_idx =
rope_3d ? emb_idx + ori_bi * 2 * max_model_len * head_dim : emb_idx;

if (hi < q_num_head + kv_num_head) {
Load<float, VecSize>(&cos_emb[new_emb_idx], &cos_emb_vec);
Expand Down Expand Up @@ -126,6 +128,7 @@ void gqa_rotary_qk_split_variable_qwen3(T *qkv_out,
const int kv_num_heads,
const int max_model_len,
const int head_dim,
const bool rope_3d,
const cudaStream_t &stream) {
assert(head_dim == 128 && "head_dim must be 128");

Expand Down Expand Up @@ -163,5 +166,6 @@ void gqa_rotary_qk_split_variable_qwen3(T *qkv_out,
num_heads,
kv_num_heads,
max_model_len,
head_dim);
head_dim,
rope_3d);
}
Loading