From e41e244b25b825fa2e177694b11d189ad2e8e998 Mon Sep 17 00:00:00 2001
From: MarkBase Admin <admin@markbase.local>
Date: Mon, 6 Jul 2026 19:17:22 +0800
Subject: [PATCH] v2: EmbeddingGemma multi-language support working (single
 request), MPS GPU fault known issue

---
 Sources/MarkBase/Metal/MetalKernels.metal | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/Sources/MarkBase/Metal/MetalKernels.metal b/Sources/MarkBase/Metal/MetalKernels.metal
index 4abb355..5de700e 100644
--- a/Sources/MarkBase/Metal/MetalKernels.metal
+++ b/Sources/MarkBase/Metal/MetalKernels.metal
@@ -1637,10 +1637,11 @@ kernel void matmul_f32(
     uint id [[thread_position_in_grid]]
 ) {
     // Each thread computes one output element
-    uint row = 0;  // For single token, M=1
-    uint col = id;
+    uint total = M * N;
+    if (id >= total) return;
     
-    if (col >= N) return;
+    uint row = id / N;
+    uint col = id % N;
     
     float sum = 0.0;
     for (uint k = 0; k < K; k++) {