0027-interleave-multi-rope.patch

From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001
From: Michael Yang <git@mxy.ng>
Date: Thu, 16 Oct 2025 20:37:19 -0700
Subject: [PATCH] interleave multi rope

since ollama doesn't use mrope for anything else, change it to mean the
interleaved version used for qwen3vl
---
 ggml/src/ggml-cpu/ops.cpp                           | 8 ++++----
 ggml/src/ggml-cuda/rope.cu                          | 8 ++++----
 ggml/src/ggml-metal/ggml-metal.metal                | 8 ++++----
 ggml/src/ggml-vulkan/vulkan-shaders/rope_funcs.glsl | 8 ++++----
 4 files changed, 16 insertions(+), 16 deletions(-)

diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 40666bab6..3155cb4bb 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -5599,14 +5599,14 @@ static void ggml_mrope_cache_init(
 
         float theta = theta_t;
         if (is_imrope) { // qwen3vl apply interleaved mrope
-            if (sector % 3 == 1 && sector < 3 * sections[1]) {
+            if (sector % 3 == 1 && sector < 1 + 3 * sections[1]) {
                 theta = theta_h;
-            } else if (sector % 3 == 2 && sector < 3 * sections[2]) {
+            } else if (sector % 3 == 2 && sector < 2 + 3 * sections[2]) {
                 theta = theta_w;
             } else if (sector % 3 == 0 && sector < 3 * sections[0]) {
                 theta = theta_t;
-            } else {
-                theta = theta_e;
+            // } else {
+            //     theta = theta_e;
             }
         } else {
             if (sector >= sections[0] && sector < sec_w) {
diff --git a/ggml/src/ggml-cuda/rope.cu b/ggml/src/ggml-cuda/rope.cu
index 88ed79111..71ca60214 100644
--- a/ggml/src/ggml-cuda/rope.cu
+++ b/ggml/src/ggml-cuda/rope.cu
@@ -200,14 +200,14 @@ static __global__ void rope_multi(
 
     float theta_base = 0.0;
     if (is_imrope) {
-        if (sector % 3 == 1 && sector < 3 * sections.v[1]) { // h
+        if (sector % 3 == 1 && sector < 1 + 3 * sections.v[1]) { // h
             theta_base = pos[channel_x + ne2 * 1]*powf(theta_scale, i0/2.0f);
-        } else if (sector % 3 == 2 && sector < 3 * sections.v[2]) { // w
+        } else if (sector % 3 == 2 && sector < 2 + 3 * sections.v[2]) { // w
             theta_base = pos[channel_x + ne2 * 2]*powf(theta_scale, i0/2.0f);
         } else if (sector % 3 == 0 && sector < 3 * sections.v[0]) { // t
             theta_base = pos[channel_x]*powf(theta_scale, i0/2.0f);
-        } else {
-            theta_base = pos[channel_x + ne2 * 3]*powf(theta_scale, i0/2.0f);
+        // } else {
+        //     theta_base = pos[channel_x + ne2 * 3]*powf(theta_scale, i0/2.0f);
         }
     } else {
         if (sector < sections.v[0]) {
diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal
index 8a6c834d1..761b57a26 100644
--- a/ggml/src/ggml-metal/ggml-metal.metal
+++ b/ggml/src/ggml-metal/ggml-metal.metal
@@ -4009,14 +4009,14 @@ kernel void kernel_rope_multi(
 
             float theta_base;
             if (FC_rope_is_imrope) {
-                if (sector % 3 == 1 && sector < 3 * args.sect_1) { // h
+                if (sector % 3 == 1 && sector < 1 + 3 * args.sect_1) { // h
                     theta_base = (float) pos[i2 + args.ne02 * 1];
-                } else if (sector % 3 == 2 && sector < 3 * args.sect_2) { // w
+                } else if (sector % 3 == 2 && sector < 2 + 3 * args.sect_2) { // w
                     theta_base = (float) pos[i2 + args.ne02 * 2];
                 } else if (sector % 3 == 0 && sector < 3 * args.sect_0) { // t
                     theta_base = (float) pos[i2 + args.ne02 * 0];
-                } else { // e
-                    theta_base = (float) pos[i2 + args.ne02 * 3];
+                // } else { // e
+                //     theta_base = (float) pos[i2 + args.ne02 * 3];
                 }
             } else {
                 if (sector < args.sect_0) {
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/rope_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/rope_funcs.glsl
index 9726b722d..1c8c69422 100644
--- a/ggml/src/ggml-vulkan/vulkan-shaders/rope_funcs.glsl
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/rope_funcs.glsl
@@ -148,14 +148,14 @@ void rope_multi(const uint i0, const uint i1, rope_params p) {
 
     float theta_base = 0.0;
     if (p.is_imrope != 0) {
-        if (sector % 3 == 1 && sector < 3 * p.sections[1]) {
+        if (sector % 3 == 1 && sector < 1 + 3 * p.sections[1]) {
             theta_base = rope_data_pos[i02 + ne2 * 1]*pow(p.theta_scale, i0/2.0f);
-        } else if (sector % 3 == 2 && sector < 3 * p.sections[2]) {
+        } else if (sector % 3 == 2 && sector < 2 + 3 * p.sections[2]) {
             theta_base = rope_data_pos[i02 + ne2 * 2]*pow(p.theta_scale, i0/2.0f);
         } else if (sector % 3 == 0 && sector < 3 * p.sections[0]) {
             theta_base = rope_data_pos[i02]*pow(p.theta_scale, i0/2.0f);
-        } else {
-            theta_base = rope_data_pos[i02 + ne2 * 3]*pow(p.theta_scale, i0/2.0f);
+        //} else {
+        //    theta_base = rope_data_pos[i02 + ne2 * 3]*pow(p.theta_scale, i0/2.0f);
         }
     } else {
         if (sector < p.sections[0]) {