diff --git a/src/liger_kernel/ops/llama4_rope.py b/src/liger_kernel/ops/llama4_rope.py index 9167d69f3..02a8647be 100644 --- a/src/liger_kernel/ops/llama4_rope.py +++ b/src/liger_kernel/ops/llama4_rope.py @@ -39,7 +39,7 @@ def _llama4_rope_kernel( Grid: (batch*seq, head) """ # 2D grid - pid_bs = tl.program_id(0) # over batch*seq + pid_bs = tl.program_id(0).to(tl.int64) # over batch*seq pid_h = tl.program_id(1) # over heads batch_idx = pid_bs // seq_len diff --git a/src/liger_kernel/ops/qwen2vl_mrope.py b/src/liger_kernel/ops/qwen2vl_mrope.py index fbd120f96..b431c8976 100644 --- a/src/liger_kernel/ops/qwen2vl_mrope.py +++ b/src/liger_kernel/ops/qwen2vl_mrope.py @@ -22,7 +22,7 @@ def _triton_qwen2vl_mrope( BLOCK_SIZE: tl.constexpr, BACKWARD_PASS: tl.constexpr = False, ): - pid = tl.program_id(0) + pid = tl.program_id(0).to(tl.int64) # locate start address q_ptr = q_ptr + pid * (n_qh * hd)