[QnA]: Why `cos` and `sin` is expected to be `hdim`, not `hdim//2`?

[tjtanaa]

I have a question regarding to the qwen2_vl MRope. From my understanding is as follows:

full_cos = torch.cat([cos_halfdim, cos_halfdim], dim=-1)
full_sin = torch.cat([sin_halfdim, sin_halfdim], dim=-1)

However from the unit tests and the code here,

Liger-Kernel/src/liger_kernel/ops/qwen2vl_mrope.py

Line 7 in 2845fe8

def _triton_qwen2vl_mrope(

the cos and sin are full_cos and full_sin instead. Is there a reason to not just pass half of the cos and sin to save memory movement?

[Tcc0403]

I believe it only reads first half of cos and sin in triton code.

Liger-Kernel/src/liger_kernel/ops/qwen2vl_mrope.py

Lines 40 to 41 in 2845fe8

	# 2. We only need the left half of cos and sin matrix because the right half is just
	# a clone of the left half.

Liger-Kernel/src/liger_kernel/ops/qwen2vl_mrope.py

Lines 52 to 61 in 2845fe8

	cos_offsets = tl.arange(0, pad_hd // 2)
	t_mask = cos_offsets < t_end
	h_mask = (t_end <= cos_offsets) & (cos_offsets < h_end)
	w_mask = (h_end <= cos_offsets) & (cos_offsets < hd // 2)
	t_cos_row = tl.load(t_cos + cos_offsets, mask=t_mask, other=0)
	h_cos_row = tl.load(h_cos + cos_offsets, mask=h_mask, other=0)
	w_cos_row = tl.load(w_cos + cos_offsets, mask=w_mask, other=0)
	t_sin_row = tl.load(t_sin + cos_offsets, mask=t_mask, other=0)
	h_sin_row = tl.load(h_sin + cos_offsets, mask=h_mask, other=0)
	w_sin_row = tl.load(w_sin + cos_offsets, mask=w_mask, other=0)

[tjtanaa]

@Tcc0403 I am getting segfault if I pass in cos_halfdim and sin_halfdim where sin_halfdim

Liger-Kernel/src/liger_kernel/ops/qwen2vl_mrope.py

Lines 45 to 50 in 2845fe8

	t_cos = cos + pid * hd
	h_cos = t_cos + bs * sl * hd
	w_cos = h_cos + bs * sl * hd
	t_sin = sin + pid * hd
	h_sin = t_sin + bs * sl * hd
	w_sin = h_sin + bs * sl * hd

When these are computed, it is using full hd to compute the pointer. So passing in cos_halfdim and sin_halfdim will cause segfault

[Tcc0403]

Oh yes, it will cause segfault if you only pass halfdim. I was saying the second half of the cos and sin aren't involved in memory movement between global and on-chip memory.

https://github.com/huggingface/transformers/blob/06f8004e5cd9d06cfbffc3f47afb6c2b43bcb3d2/src/transformers/models/qwen2_vl/modeling_qwen2_vl.py#L178

For example, note
that cos[position_ids] and sin[position_ids] have the shape [batch_size, seq_len, head_dim].

I think liger rope was designed to match the api (assuming both are full shape), and only reading the first half in kernel code.

If you want to save some memory allocation by passing half cos and sin only, I think we can make an additional check and create another parameter such as cos_hd to avoid segfault on indexing for those tensors.