server: cap shared decision prefixes at the first output (#30710)

Limit the shared prefix to the first required output position of the parent and each child task.
This keeps decision output embeddings available in the final batch when questions share option tokens.

Fixes #30482

Signed-off-by: aiqubits <aiqubits@hotmail.com>

common/arg.cpp

@@ -12,16 +12,6 @@
 #include "speculative.h"
 #include "preset.h"
 
-// fix problem with std::min and std::max
-#if defined(_WIN32)
-#define WIN32_LEAN_AND_MEAN
-#ifndef NOMINMAX
-#   define NOMINMAX
-#endif
-#include <windows.h>
-#include <shellapi.h>
-#endif
-
 #include <algorithm>
 #include <cinttypes>
 #include <climits>
@@ -38,20 +28,6 @@
 #include <thread> // for hardware_concurrency
 #include <vector>
 
-#ifndef __EMSCRIPTEN__
-#ifdef __linux__
-#include <linux/limits.h>
-#elif defined(_WIN32)
-#   if !defined(PATH_MAX)
-#   define PATH_MAX MAX_PATH
-#   endif
-#elif defined(_AIX)
-#include <sys/limits.h>
-#else
-#include <sys/syslimits.h>
-#endif
-#endif
-
 #define LLAMA_MAX_URL_LENGTH 2084 // Maximum URL Length in Chrome: 2083
 
 using json = common_json;
@@ -351,7 +327,7 @@ static bool parse_bool_value(const std::string & value) {
 static std::string get_default_local_path(const std::string & url) {
     auto f = string_split<std::string>(url, '#').front();
     f = string_split<std::string>(f, '?').front();
-    return fs_path_to_utf8(fs_get_cache_file(string_split<std::string>(f, '/').back()));
+    return fs_get_cache_file(string_split<std::string>(f, '/').back()).string();
 }
 
 static bool spec_types_is_default(const common_params & params) {
@@ -762,7 +738,7 @@ static void common_params_apply_system_config(common_params & params, llama_exam
     common_preset_context ctx(ex);
     ctx.ignore_unknown_keys = true; // the same config file is shared by all programs
     for (const auto & path : found) {
-        LOG_INF("using config file: %s\n", fs_path_to_utf8(path).c_str());
+        LOG_INF("using config file: %s\n", path.string().c_str());
         common_preset global;
         common_presets presets = ctx.load_from_ini(path, global);
         global.apply_to_params(params);
@@ -1259,44 +1235,7 @@ bool common_params_to_map(int argc, char ** argv, llama_example ex, std::map<com
     return true;
 }
 
-#ifdef _WIN32
-struct utf8_argv {
-    std::vector<std::string> buf;
-    std::vector<char*> ptrs;
-};
-
-static utf8_argv make_utf8_argv() {
-    utf8_argv out;
-    int wargc = 0;
-    LPWSTR* wargv = CommandLineToArgvW(GetCommandLineW(), &wargc);
-    if (!wargv) return out;
-
-    out.buf.reserve(wargc);
-    for (int i = 0; i < wargc; ++i) {
-        int n = WideCharToMultiByte(CP_UTF8, WC_ERR_INVALID_CHARS, wargv[i], -1, nullptr, 0, nullptr, nullptr);
-        if (n <= 0) { out.buf.emplace_back(); continue; }
-        auto& s = out.buf.emplace_back();
-        s.resize(static_cast<size_t>(n - 1));
-        (void)WideCharToMultiByte(CP_UTF8, 0, wargv[i], -1, s.data(), n, nullptr, nullptr);
-    }
-    LocalFree(wargv);
-
-    out.ptrs.reserve(out.buf.size() + 1);
-    for (auto& s : out.buf) out.ptrs.push_back(s.data());
-    out.ptrs.push_back(nullptr);
-    return out;
-}
-#endif
-
 bool common_params_parse(int argc, char ** argv, common_params & params, llama_example ex, void(*print_usage)(int, char **)) {
-#ifdef _WIN32
-    auto utf8 = make_utf8_argv();
-    // repair argv only when it matches the process command line
-    if (static_cast<int>(utf8.buf.size()) == argc) {
-        argv = utf8.ptrs.data();
-    }
-#endif
-
     auto ctx_arg = common_params_parser_init(params, ex, print_usage);
     const common_params params_org = ctx_arg.params; // the example can modify the default params
 

common/common.cpp

@@ -934,12 +934,6 @@ std::string wstring_to_utf8(const std::wstring & str) {
 }
 #endif
 
-// returns the path as a UTF-8 string, preserving its separators
-std::string fs_path_to_utf8(const std::filesystem::path & path) {
-    const auto value = path.u8string();
-    return std::string(value.begin(), value.end());
-}
-
 void fs_write_atomic(const std::filesystem::path & path, const std::string & data) {
     // a random suffix gives every writer its own temporary file, across processes and containers
     std::random_device rd;
@@ -963,7 +957,7 @@ void fs_write_atomic(const std::filesystem::path & path, const std::string & dat
 
     if (file.fail() || ec) {
         std::filesystem::remove(path_tmp, ec);
-        throw std::runtime_error("failed to write file: " + fs_path_to_utf8(path));
+        throw std::runtime_error("failed to write file: " + path.string());
     }
 }
 
@@ -1057,7 +1051,7 @@ std::filesystem::path fs_get_cache_file(const std::string & filename) {
     std::error_code ec;
     common_create_directories(cache_directory, ec);
     if (ec) {
-        throw std::runtime_error("failed to create cache directory: " + fs_path_to_utf8(cache_directory));
+        throw std::runtime_error("failed to create cache directory: " + cache_directory.string());
     }
     return cache_directory / std::filesystem::u8path(filename);
 }

common/common.h

@@ -910,9 +910,6 @@ std::wstring utf8_to_wstring(const std::string & str);
 std::string  wstring_to_utf8(const std::wstring & str);
 #endif
 
-// returns the path as a UTF-8 string, preserving its separators
-std::string fs_path_to_utf8(const std::filesystem::path & path);
-
 //
 // Environment utils
 //

common/download.cpp

@@ -1005,7 +1005,7 @@ std::string common_docker_resolve_model(const std::string & docker) {
         std::string model_filename = repo;
         std::replace(model_filename.begin(), model_filename.end(), '/', '_');
         model_filename += "_" + tag + ".gguf";
-        std::string local_path = fs_path_to_utf8(fs_get_cache_file(model_filename));
+        std::string local_path = fs_get_cache_file(model_filename).string();
 
         const std::string blob_url = url_prefix + "/blobs/" + gguf_digest;
         common_download_opts opts;

common/hf-cache.cpp

@@ -63,7 +63,7 @@ static fs::path get_cache_directory() {
 }
 
 std::string get_cache_path() {
-    return fs_path_to_utf8(get_cache_directory());
+    return get_cache_directory().string();
 }
 
 static std::string folder_name_to_repo(const std::string & folder) {
@@ -340,11 +340,11 @@ hf_files get_repo_files(const std::string & repo_id,
             file.url = endpoint + repo_id + "/resolve/" + commit + "/" + file.path;
 
             fs::path final_path = commit_path / subpath;
-            file.final_path = fs_path_to_utf8(final_path);
+            file.final_path = final_path.string();
 
             if (!file.oid.empty() && !fs::exists(final_path)) {
                 fs::path local_path = blobs_path / file.oid;
-                file.local_path = fs_path_to_utf8(local_path);
+                file.local_path = local_path.string();
             } else {
                 file.local_path = file.final_path;
             }
@@ -411,7 +411,7 @@ hf_files get_cached_files(const std::string & repo_id) {
         if (!fs::exists(snapshots_path)) {
             continue;
         }
-        std::string _repo_id = folder_name_to_repo(fs_path_to_utf8(repo.path().filename()));
+        std::string _repo_id = folder_name_to_repo(repo.path().filename().string());
 
         if (!is_valid_repo_id(_repo_id)) {
             continue;
@@ -436,7 +436,7 @@ hf_files get_cached_files(const std::string & repo_id) {
                 file.repo_id = _repo_id;
                 const auto generic_path = path.generic_u8string();
                 file.path = std::string(generic_path.begin(), generic_path.end());
-                file.local_path = fs_path_to_utf8(entry.path());
+                file.local_path = entry.path().string();
                 file.final_path = file.local_path;
                 files.push_back(std::move(file));
             }
@@ -516,7 +516,7 @@ bool remove_cached_repo(const std::string & repo_id) {
     std::error_code ec;
     auto removed = fs::remove_all(repo_path, ec);
     if (ec) {
-        LOG_ERR("%s: failed to remove repo cache %s: %s\n", __func__, fs_path_to_utf8(repo_path).c_str(), ec.message().c_str());
+        LOG_ERR("%s: failed to remove repo cache %s: %s\n", __func__, repo_path.string().c_str(), ec.message().c_str());
         return false;
     }
     return removed > 0;

common/parsers/qwen3-coder.cpp

@@ -74,16 +74,24 @@ common_chat_params common_chat_params_init_qwen3_coder(const common_chat_templat
     data.parser = build_chat_peg_parser([&](common_chat_peg_builder & p) {
         auto generation_prompt = p.literal(GEN_PREFIX);
 
+        // The template strips the newline before </think>, so leave it out of the reasoning
+        auto think_body  = p.until_one_of({ "\n</think>", "</think>", "<tool_call>" });
+        auto think_close = p.literal("\n</think>") | p.literal("</think>") | p.peek(p.literal("<tool_call>"));
+
+        auto reasoning_extracted = p.optional("<think>" + p.space() + p.reasoning(think_body) + think_close);
+        auto reasoning_skipped   = p.optional("<think>" + p.space() + think_body + think_close);
+
         auto reasoning = p.eps();
         if (supports_reasoning && extract_reasoning) {
-            // The template strips the newline before </think>, so leave it out of the reasoning
-            reasoning = p.optional("<think>" + p.space() +
-                                   p.reasoning(p.until_one_of({ "\n</think>", "</think>", "<tool_call>" })) +
-                                   (p.literal("\n</think>") | p.literal("</think>") | p.peek(p.literal("<tool_call>"))));
+            reasoning = reasoning_extracted;
         }
 
         // Response format parser
         if (has_response_format) {
+            // The answer must be bare JSON, so the think block is consumed even when it is not extracted
+            if (supports_reasoning && !extract_reasoning) {
+                reasoning = reasoning_skipped;
+            }
             return generation_prompt + (reasoning << p.content(p.schema(p.json(), "response-format", inputs.json_schema)));
         }
 

common/preset.cpp

@@ -171,12 +171,12 @@ static std::map<std::string, std::map<std::string, std::string>> parse_ini_from_
     std::map<std::string, std::map<std::string, std::string>> parsed;
 
     if (!std::filesystem::exists(path)) {
-        throw std::runtime_error("preset file does not exist: " + fs_path_to_utf8(path));
+        throw std::runtime_error("preset file does not exist: " + path.string());
     }
 
     std::ifstream file(path);
     if (!file.good()) {
-        throw std::runtime_error("failed to open server preset file: " + fs_path_to_utf8(path));
+        throw std::runtime_error("failed to open server preset file: " + path.string());
     }
 
     std::string contents((std::istreambuf_iterator<char>(file)), std::istreambuf_iterator<char>());
@@ -225,7 +225,7 @@ static std::map<std::string, std::map<std::string, std::string>> parse_ini_from_
     common_peg_parse_context ctx(contents);
     const auto result = parser.parse(ctx);
     if (!result.success()) {
-        throw std::runtime_error("failed to parse server config file: " + fs_path_to_utf8(path));
+        throw std::runtime_error("failed to parse server config file: " + path.string());
     }
 
     std::string current_section = COMMON_PRESET_DEFAULT_NAME;
@@ -323,7 +323,7 @@ common_presets common_preset_context::load_from_ini(const std::filesystem::path
                 }
                 LOG_DBG("accepted option: %s = %s\n", key.c_str(), preset.options[opt].c_str());
             } else if (ignore_unknown_keys) {
-                LOG_WRN("ignoring option '%s' from %s: not supported by this program\n", key.c_str(), fs_path_to_utf8(path).c_str());
+                LOG_WRN("ignoring option '%s' from %s: not supported by this program\n", key.c_str(), path.string().c_str());
             } else {
                 throw std::runtime_error(string_format(
                     "option '%s' not recognized in preset '%s'",
@@ -401,7 +401,7 @@ common_presets common_preset_context::load_from_models_dir(const std::string & m
             if (!entry.is_regular_file(ec)) {
                 continue;
             }
-            const std::string fname = fs_path_to_utf8(entry.path().filename());
+            const std::string fname = entry.path().filename().string();
             if (!string_ends_with(fname, ".gguf")) {
                 continue;
             }
@@ -422,9 +422,9 @@ common_presets common_preset_context::load_from_models_dir(const std::string & m
         if (!path.empty()) {
             models.push_back({
                 /* name        */ name,
-                /* path        */ fs_path_to_utf8(path),
-                /* path_mmproj */ fs_path_to_utf8(mmproj_file), // can be empty
-                /* path_draft  */ fs_path_to_utf8(draft_file)    // can be empty
+                /* path        */ path.string(),
+                /* path_mmproj */ mmproj_file.string(), // can be empty
+                /* path_draft  */ draft_file.string()   // can be empty
             });
         }
     };
@@ -432,13 +432,13 @@ common_presets common_preset_context::load_from_models_dir(const std::string & m
     for (const auto & entry : std::filesystem::directory_iterator(dir)) {
         std::error_code ec;
         if (entry.is_directory(ec)) {
-            scan_subdir(entry.path(), fs_path_to_utf8(entry.path().filename()));
+            scan_subdir(entry.path(), entry.path().filename().string());
             continue;
         }
         if (!entry.is_regular_file(ec)) {
             continue;
         }
-        const std::string fname = fs_path_to_utf8(entry.path().filename());
+        const std::string fname = entry.path().filename().string();
         if (!string_ends_with(fname, ".gguf")) {
             continue;
         }
@@ -450,7 +450,7 @@ common_presets common_preset_context::load_from_models_dir(const std::string & m
         string_replace_all(name, ".gguf", "");
         models.push_back({
             /* name        */ name,
-            /* path        */ fs_path_to_utf8(entry.path()),
+            /* path        */ entry.path().string(),
             /* path_mmproj */ "",
             /* path_draft  */ ""
         });

docs/ops.md

@@ -74,7 +74,7 @@ Legend:
 |                             MEAN | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | 🟡 | ✅ | ❌ | ❌ | ❌ |
 |                              MUL | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                          MUL_MAT | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 |
-|                 MUL_MAT_HADAMARD | ❌ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ✅ | ❌ | 🟡 |
+|                 MUL_MAT_HADAMARD | ❌ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | 🟡 |
 |                       MUL_MAT_ID | ❌ | 🟡 | ✅ | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | 🟡 | ❌ |
 |                  MUL_MAT_ID_W4A4 | ❌ | ❌ | ✅ | ❌ | ❌ | 🟡 | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ | ❌ |
 |                  MUL_MAT_ID_W4A8 | ❌ | ❌ | ✅ | ❌ | ❌ | 🟡 | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ | ❌ |

docs/ops/OpenCL.csv

@@ -9242,31 +9242,31 @@
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=128,n=4,k=128,bs=[2,3],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=32,n=1,k=32,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=16384,n=1,k=16384,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=4,k=128,bs=[2,3],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=4,k=128,bs=[2,3],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=1024,n=1,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=2048,n=1,k=2048,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=4096,n=1,k=4096,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=8192,n=1,k=8192,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=1024,n=7,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=1,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=2048,n=1,k=2048,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=4096,n=1,k=4096,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=8192,n=1,k=8192,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=7,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=1,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=2048,n=1,k=2048,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=4096,n=1,k=4096,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=8192,n=1,k=8192,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=7,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
 "OpenCL","MUL_MAT_W4A8","type_a=nvfp4,type_b=f32,m=32,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
 "OpenCL","MUL_MAT_W4A8","type_a=nvfp4,type_b=f32,m=32,n=32,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
 "OpenCL","MUL_MAT_W4A8","type_a=nvfp4,type_b=f32,m=64,n=16,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"

ggml/src/ggml-metal/kernels/quantize.h

@@ -20,6 +20,7 @@ void quantize_q1_0(device const float * src, device block_q1_0 & dst) {
 }
 
 void quantize_q2_0(device const float * src, device block_q2_0 & dst) {
+#pragma METAL fp math_mode(safe)
     float amax = 0.0f;
     for (int j = 0; j < QK2_0; j++) {
         float a = fabs(src[j]);

ggml/src/ggml-opencl/CMakeLists.txt

@@ -70,6 +70,7 @@ set(GGML_OPENCL_KERNELS
     lightning_indexer
     tri
     fill
+    fwht
     clamp
     cpy
     cvt
@@ -117,6 +118,12 @@ set(GGML_OPENCL_KERNELS
     gemm_noshuffle_q2_0_f32
     gemm_noshuffle_q2_0_q8_1_dp4a
     gemv_noshuffle_q2_0_f32
+    gemm_noshuffle_iq1_m_f32
+    gemm_noshuffle_iq1_m_q8_1_dp4a
+    gemv_noshuffle_iq1_m_f32
+    gemm_noshuffle_iq1_s_f32
+    gemm_noshuffle_iq1_s_q8_1_dp4a
+    gemv_noshuffle_iq1_s_f32
     gemm_noshuffle_q3_k_q8_1_dp4a
     gemv_noshuffle_q3_k_f32
     mul_mv_q2_k_f32
@@ -136,6 +143,13 @@ set(GGML_OPENCL_KERNELS
     mul_mv_q2_0_f32
     mul_mv_q2_0_f32_flat
     mul_mm_q2_0_f32_l4_lm
+    mul_mv_iq1_m_f32
+    mul_mv_iq1_m_f32_flat
+    mul_mm_iq1_m_f32_l4_lm
+    iq1_s_grid
+    mul_mv_iq1_s_f32
+    mul_mv_iq1_s_f32_flat
+    mul_mm_iq1_s_f32_l4_lm
     mul_mv_q3_k_f32_flat
     mul_mm_q3_k_f32_l4_lm
     mul_mv_q2_k_f32_flat

ggml/src/ggml-opencl/kernels/cvt.cl

@@ -3088,3 +3088,93 @@ kernel void kernel_restore_block_q2_0(
         dst[i].qs[j] = src_qs[16*i + j];
     }
 }
+
+struct block_iq1_s {
+    half d;
+    uchar qs[QK_K/8];
+    ushort qh[QK_K/32];
+};
+
+kernel void kernel_convert_block_iq1_s(
+    global struct block_iq1_s * src0,
+    global uchar * dst_qs,
+    global ushort * dst_qh,
+    global half * dst_d,
+    ulong n_blocks
+) {
+    const ulong i = get_global_id(0);
+    if (i >= n_blocks) { return; }
+    dst_d[i] = src0[i].d;
+    for (int j = 0; j < 32; ++j) {
+        dst_qs[32*i + j] = src0[i].qs[j];
+    }
+    for (int j = 0; j < 8; ++j) {
+        dst_qh[8*i + j] = src0[i].qh[j];
+    }
+}
+
+kernel void kernel_restore_block_iq1_s(
+    global uchar * src_qs,
+    global ushort * src_qh,
+    global half * src_d,
+    global struct block_iq1_s * dst,
+    ulong n_blocks
+) {
+    const ulong i = get_global_id(0);
+    if (i >= n_blocks) { return; }
+    dst[i].d = src_d[i];
+    for (int j = 0; j < 32; ++j) {
+        dst[i].qs[j] = src_qs[32*i + j];
+    }
+    for (int j = 0; j < 8; ++j) {
+        dst[i].qh[j] = src_qh[8*i + j];
+    }
+}
+
+struct block_iq1_m {
+    uchar qs[QK_K/8];
+    uchar qh[QK_K/16];
+    uchar scales[QK_K/32];
+};
+
+kernel void kernel_convert_block_iq1_m(
+    global struct block_iq1_m * src0,
+    global uchar * dst_qs,
+    global uchar * dst_qh,
+    global ushort * dst_sc,
+    ulong n_blocks
+) {
+    const ulong i = get_global_id(0);
+    if (i >= n_blocks) { return; }
+    for (int j = 0; j < 32; ++j) {
+        dst_qs[32*i + j] = src0[i].qs[j];
+    }
+    for (int j = 0; j < 16; ++j) {
+        dst_qh[16*i + j] = src0[i].qh[j];
+    }
+    for (int j = 0; j < 4; ++j) {
+        dst_sc[4*i + j] = (ushort)src0[i].scales[2*j] | ((ushort)src0[i].scales[2*j + 1] << 8);
+    }
+}
+
+kernel void kernel_restore_block_iq1_m(
+    global uchar * src_qs,
+    global uchar * src_qh,
+    global ushort * src_sc,
+    global struct block_iq1_m * dst,
+    ulong n_blocks
+) {
+    const ulong i = get_global_id(0);
+    if (i >= n_blocks) { return; }
+    for (int j = 0; j < 32; ++j) {
+        dst[i].qs[j] = src_qs[32*i + j];
+    }
+    for (int j = 0; j < 16; ++j) {
+        dst[i].qh[j] = src_qh[16*i + j];
+    }
+    for (int j = 0; j < 4; ++j) {
+        ushort sc = src_sc[4*i + j];
+        dst[i].scales[2*j] = (uchar)sc;
+        dst[i].scales[2*j + 1] = (uchar)(sc >> 8);
+    }
+}

ggml/src/ggml-opencl/kernels/fwht.cl

@@ -0,0 +1,44 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_fwht(
+        global const char * src,
+        ulong offset_src,
+        global char * dst,
+        ulong offset_dst,
+        int n,
+        ulong ne1,
+        ulong ne2,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3,
+        float scale,
+        local float * work) {
+    const ulong row = get_group_id(1);
+    const ulong i1 = row % ne1;
+    const ulong i2 = row / ne1 % ne2;
+    const ulong i3 = row / ne1 / ne2;
+    global const SRC_TYPE * x = (global const SRC_TYPE *)(src + offset_src + i1*nb1 + i2*nb2 + i3*nb3);
+    global float * y = (global float *)(dst + offset_dst) + row*n;
+    const int tid = get_local_id(0);
+    const int nth = get_local_size(0);
+
+    for (int i = tid; i < n; i += nth) {
+        work[i] = (float)x[i] * scale;
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    for (int h = 1; h < n; h *= 2) {
+        for (int i = tid; i < n/2; i += nth) {
+            const int j = 2*i - (i & (h - 1));
+            const float a = work[j];
+            const float b = work[j + h];
+            work[j] = a + b;
+            work[j + h] = a - b;
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    for (int i = tid; i < n; i += nth) {
+        y[i] = work[i];
+    }
+}

ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq1_m_f32.cl

@@ -0,0 +1,93 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_qcom_reqd_sub_group_size
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+__attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+kernel void kernel_gemm_noshuffle_iq1_m_f32(
+    global const uchar * src0_qs,
+    global const uchar * src0_qh,
+    global const ushort * src0_sc,
+    read_only image1d_buffer_t src1,
+    global float * dst,
+    ulong offsetd,
+    int m,
+    int n,
+    int k,
+    int n_no_padding
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+    const int gy = get_global_id(0);
+    const int row = get_global_id(1) * 4;
+    const int n_4 = n / 4;
+    if (row >= m) {
+        return;
+    }
+
+    half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
+    for (int i = 0; i < k; i += 16) {
+        const uint4 w0 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 0) * m));
+        const uint4 w1 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 1) * m));
+        const uint4 w2 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 2) * m));
+        const uint4 w3 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 3) * m));
+        const ushort4 bits = convert_ushort4((w0 >> (uint4)(12)) | ((w1 >> (uint4)(8)) & (uint4)(0x00f0))
+                                           | ((w2 >> (uint4)(4)) & (uint4)(0x0f00)) | (w3 & (uint4)(0xf000)));
+        const half4 d = as_half4(bits);
+        const uint4 sc = convert_uint4(vload4(0, src0_sc + row + (i / 64) * m));
+        const half4 scale = d * convert_half4(2*((sc >> (uint4)(3*((i / 16) & 3))) & (uint4)(7)) + 1);
+        const uint4 qh = convert_uint4(vload4(0, src0_qh + row + (i / 16) * m));
+
+        for (int l = 0; l < 16; l += 4) {
+            const int ki = i + l;
+            const uint4 gi = convert_uint4(vload4(0, src0_qs + row + (ki / 8) * m))
+                           | (((qh >> (uint4)(4*((ki / 8) & 1))) & (uint4)(7)) << (uint4)(8));
+            const uint4 grid = (uint4)(iq1_s_grid[gi.s0], iq1_s_grid[gi.s1], iq1_s_grid[gi.s2], iq1_s_grid[gi.s3]);
+            const float4 delta = select((float4)(-0.875f), (float4)(-1.125f), (qh & (uint4)(0x08 << (4*((ki / 8) & 1)))) != 0);
+            #pragma unroll
+            for (int j = 0; j < 4; ++j) {
+                half8 b;
+                b.s0123 = read_imageh(src1, gy * 2     + (ki + j) * n_4);
+                b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4);
+                const float4 magnitude = convert_float4((grid >> (uint4)(8*j + (ki & 4))) & (uint4)(15));
+                const half4 w = convert_half4(magnitude + delta) * scale;
+                c0 += b * w.s0;
+                c1 += b * w.s1;
+                c2 += b * w.s2;
+                c3 += b * w.s3;
+            }
+        }
+    }
+
+    int idx = gy * 8 * m + row;
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
+    }
+}

ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq1_m_q8_1_dp4a.cl

@@ -0,0 +1,144 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#ifdef cl_khr_integer_dot_product
+#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
+#endif
+
+#define QK_K 256
+
+#ifndef TILESIZE_N
+#define TILESIZE_N 32
+#endif
+
+inline uint iq1_m_pack(uint g, uint shift, uint negative_delta) {
+    // 8*(grid + delta) is an exact signed integer in [-9, 9].
+    const int bias = negative_delta ? -9 : -7;
+    const int4 values = 8*convert_int4(as_uchar4((g >> shift) & 0x0F0F0F0Fu)) + bias;
+    return as_uint(convert_char4(values));
+}
+
+// Pass staged operands by value; local dp4a operands miscompile on X2.
+inline int dot4_q8a_v(uint4 qw, uint4 a) {
+    int r = 0;
+    r = dot_acc_sat_4x8packed_ss_int(qw.s0, a.x, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s1, a.y, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s2, a.z, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s3, a.w, r);
+    return r;
+}
+
+__attribute__((qcom_wave_pair_mode(1)))
+kernel void kernel_gemm_noshuffle_iq1_m_q8_1_dp4a(
+        __global const uchar  * src0_qs,
+        __global const uchar  * src0_qh,
+        __global const ushort * src0_sc,
+        __global const uint   * src1_qa,
+        __global const half   * src1_da,
+        __global       float  * dst,
+        ulong  offsetd,
+        int    m,
+        int    n_no_padding,
+        int    k
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+
+    const uint lid = get_local_id(0);
+    const uint block_id_m = get_global_id(1);
+    const uint block_id_n = get_global_id(2);
+
+    const uint row      = block_id_m * 64 + lid;
+    const uint col_base = block_id_n * TILESIZE_N;
+    const bool row_valid = row < (uint)m;
+    const uint rrow     = row_valid ? row : 0;
+
+    const uint k_u = (uint)k >> 2;
+    const uint k_b = (uint)k >> 5;
+
+    __local uint4 sh_qa4[TILESIZE_N][2];
+    __local half sh_d[TILESIZE_N];
+
+#define IQ1_M_GRID(i) iq1_s_grid[(i)]
+
+#define NGROUPS (TILESIZE_N / 4)
+    float4 acc[NGROUPS];
+    #pragma unroll
+    for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f);
+
+    float dsuper = 0.f;
+    for (uint step = 0; step < (uint)k; step += 32) {
+        const uint sub = step >> 5;
+        const uint ib  = sub >> 3;
+
+        if ((sub & 7u) == 0u) {
+            const uint sb = rrow + 4u*ib*(uint)m;
+            const uint w0 = src0_sc[sb + 0u*(uint)m];
+            const uint w1 = src0_sc[sb + 1u*(uint)m];
+            const uint w2 = src0_sc[sb + 2u*(uint)m];
+            const uint w3 = src0_sc[sb + 3u*(uint)m];
+            const ushort bits = (ushort)((w0 >> 12) | ((w1 >> 8) & 0x00f0u)
+                                       | ((w2 >> 4) & 0x0f00u) | (w3 & 0xf000u));
+            dsuper = (float)as_half(bits);
+        }
+        const uint sc = src0_sc[rrow + (4u*ib + (sub & 7u)/2u)*(uint)m];
+        const uint sh = 6u*(sub & 1u);
+        const float dl0 = dsuper * (float)(2u*((sc >> sh) & 7u) + 1u) * 0.125f;
+        const float dl1 = dsuper * (float)(2u*((sc >> (sh + 3u)) & 7u) + 1u) * 0.125f;
+        const uint h0 = src0_qh[rrow + (2u*sub + 0u)*(uint)m];
+        const uint h1 = src0_qh[rrow + (2u*sub + 1u)*(uint)m];
+        const uint gb = rrow + 4u*sub*(uint)m;
+        const uint g0 = IQ1_M_GRID((uint)src0_qs[gb + 0u*(uint)m] | (( h0       & 7u) << 8));
+        const uint g1 = IQ1_M_GRID((uint)src0_qs[gb + 1u*(uint)m] | (((h0 >> 4) & 7u) << 8));
+        const uint g2 = IQ1_M_GRID((uint)src0_qs[gb + 2u*(uint)m] | (( h1       & 7u) << 8));
+        const uint g3 = IQ1_M_GRID((uint)src0_qs[gb + 3u*(uint)m] | (((h1 >> 4) & 7u) << 8));
+        const uint4 qlo = (uint4)(iq1_m_pack(g0, 0, h0 & 0x08u), iq1_m_pack(g0, 4, h0 & 0x08u),
+                                 iq1_m_pack(g1, 0, h0 & 0x80u), iq1_m_pack(g1, 4, h0 & 0x80u));
+        const uint4 qhi = (uint4)(iq1_m_pack(g2, 0, h1 & 0x08u), iq1_m_pack(g2, 4, h1 & 0x08u),
+                                 iq1_m_pack(g3, 0, h1 & 0x80u), iq1_m_pack(g3, 4, h1 & 0x80u));
+
+        for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) {
+            const uint t = idx >> 1;
+            const uint v = idx & 1;
+            const uint c = col_base + t;
+            sh_qa4[t][v] = (c < (uint)n_no_padding)
+                         ? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2))
+                         : (uint4)(0u);
+        }
+        if (lid < TILESIZE_N) {
+            const uint c = col_base + lid;
+            sh_d[lid] = (c < (uint)n_no_padding) ? src1_da[c * k_b + sub] : (half)0;
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3]))
+#define IQ1_M_COL(b) (dl0 * (float)dot4_q8a_v(qlo, (uint4)(sh_qa4[b][0])) \
+                           + dl1 * (float)dot4_q8a_v(qhi, (uint4)(sh_qa4[b][1])))
+        #pragma unroll
+        for (int g = 0; g < NGROUPS; ++g) {
+            const int b = g * 4;
+            float4 rf;
+            rf.s0 = IQ1_M_COL(b+0);  rf.s1 = IQ1_M_COL(b+1);
+            rf.s2 = IQ1_M_COL(b+2);  rf.s3 = IQ1_M_COL(b+3);
+            acc[g] += LD4(sh_d, b) * rf;
+        }
+#undef IQ1_M_COL
+#undef LD4
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    if (!row_valid) {
+        return;
+    }
+
+    #pragma unroll
+    for (int g = 0; g < NGROUPS; ++g) {
+        const uint b = (uint)(g * 4);
+        const float4 a = acc[g];
+        const uint c0 = col_base + b;
+        if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0;
+        if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1;
+        if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2;
+        if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3;
+    }
+#undef NGROUPS
+#undef IQ1_M_GRID
+}

ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq1_s_f32.cl

@@ -0,0 +1,86 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_qcom_reqd_sub_group_size
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+__attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+kernel void kernel_gemm_noshuffle_iq1_s_f32(
+    global const uchar * src0_qs,
+    global const ushort * src0_qh,
+    global const half  * src0_d,
+    read_only image1d_buffer_t src1,
+    global float * dst,
+    ulong offsetd,
+    int m,
+    int n,
+    int k,
+    int n_no_padding
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+    const int gy = get_global_id(0);
+    const int row = get_global_id(1) * 4;
+    const int n_4 = n / 4;
+    if (row >= m) {
+        return;
+    }
+
+    half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
+    for (int i = 0; i < k; i += 16) {
+        const uint4 qh = convert_uint4(vload4(0, src0_qh + row + (i / 32) * m));
+        const half4 d = vload4(0, src0_d + row + (i / 256) * m);
+        const half4 scale = d * convert_half4(2*((qh >> (uint4)(12)) & (uint4)(7)) + 1);
+        const float4 delta = select((float4)(-0.875f), (float4)(-1.125f), (qh & (uint4)(0x8000)) != 0);
+
+        for (int l = 0; l < 16; l += 4) {
+            const int ki = i + l;
+            const uint4 gi = convert_uint4(vload4(0, src0_qs + row + (ki / 8) * m))
+                           | (((qh >> (uint4)(3*((ki / 8) & 3))) & (uint4)(7)) << (uint4)(8));
+            const uint4 grid = (uint4)(iq1_s_grid[gi.s0], iq1_s_grid[gi.s1], iq1_s_grid[gi.s2], iq1_s_grid[gi.s3]);
+            #pragma unroll
+            for (int j = 0; j < 4; ++j) {
+                half8 b;
+                b.s0123 = read_imageh(src1, gy * 2     + (ki + j) * n_4);
+                b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4);
+                const float4 magnitude = convert_float4((grid >> (uint4)(8*j + (ki & 4))) & (uint4)(15));
+                const half4 w = convert_half4(magnitude + delta) * scale;
+                c0 += b * w.s0;
+                c1 += b * w.s1;
+                c2 += b * w.s2;
+                c3 += b * w.s3;
+            }
+        }
+    }
+
+    int idx = gy * 8 * m + row;
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx);
+        idx += m;
+    }
+    if (idx + 3 < m * n_no_padding) {
+        vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
+    }
+}

ggml/src/ggml-opencl/kernels/gemm_noshuffle_iq1_s_q8_1_dp4a.cl

@@ -0,0 +1,138 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#ifdef cl_khr_integer_dot_product
+#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
+#endif
+
+#define QK_K 256
+
+#ifndef TILESIZE_N
+#define TILESIZE_N 32
+#endif
+
+// Pass staged operands by value; local dp4a operands miscompile on X2.
+inline int dot8_q8a_v(uint8 qw, uint4 a0, uint4 a1) {
+    int r = 0;
+    r = dot_acc_sat_4x8packed_ss_int(qw.s0, a0.x, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s1, a0.y, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s2, a0.z, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s3, a0.w, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s4, a1.x, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s5, a1.y, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s6, a1.z, r);
+    r = dot_acc_sat_4x8packed_ss_int(qw.s7, a1.w, r);
+    return r;
+}
+
+__attribute__((qcom_wave_pair_mode(1)))
+kernel void kernel_gemm_noshuffle_iq1_s_q8_1_dp4a(
+        __global const uchar  * src0_qs,
+        __global const ushort * src0_qh,
+        __global const half   * src0_d,
+        __global const uint   * src1_qa,
+        __global const half   * src1_da,
+        __global const half   * src1_sa,
+        __global       float  * dst,
+        ulong  offsetd,
+        int    m,
+        int    n_no_padding,
+        int    k
+) {
+    dst = (global float *)((global char *)dst + offsetd);
+
+    const uint lid = get_local_id(0);
+    const uint block_id_m = get_global_id(1);
+    const uint block_id_n = get_global_id(2);
+
+    const uint row      = block_id_m * 64 + lid;
+    const uint col_base = block_id_n * TILESIZE_N;
+    const bool row_valid = row < (uint)m;
+    const uint rrow     = row_valid ? row : 0;
+
+    const uint k_u = (uint)k >> 2;
+    const uint k_b = (uint)k >> 5;
+
+    __local uint4 sh_qa4[TILESIZE_N][2];
+    __local half sh_da[TILESIZE_N];
+    __local half sh_sa[TILESIZE_N];
+
+#define IQ1_S_GRID(i) iq1_s_grid[(i)]
+
+#define NGROUPS (TILESIZE_N / 4)
+    float4 acc[NGROUPS];
+    #pragma unroll
+    for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f);
+
+    for (uint step = 0; step < (uint)k; step += 32) {
+        const uint sub = step >> 5;
+        const uint ib  = sub >> 3;
+
+        const uint  qhv = (uint)src0_qh[rrow + sub * (uint)m];
+        const float dl  = (float)src0_d[rrow + ib * (uint)m]
+                        * (float)(2u * ((qhv >> 12) & 7u) + 1u);
+
+        // The codebook stores grid + 1; sa includes the Q8_1 activation scale.
+        const float dlt = ((qhv & 0x8000u) ? -0.125f : 0.125f) - 1.0f;
+
+        const uint gb = rrow + (step >> 3) * (uint)m;
+
+        uint8 qw;
+        {
+            const uint g0 = IQ1_S_GRID((uint)src0_qs[gb + 0u * (uint)m] | ((( qhv       & 7u) << 8)));
+            const uint g1 = IQ1_S_GRID((uint)src0_qs[gb + 1u * (uint)m] | ((((qhv >> 3) & 7u) << 8)));
+            const uint g2 = IQ1_S_GRID((uint)src0_qs[gb + 2u * (uint)m] | ((((qhv >> 6) & 7u) << 8)));
+            const uint g3 = IQ1_S_GRID((uint)src0_qs[gb + 3u * (uint)m] | ((((qhv >> 9) & 7u) << 8)));
+            qw.s0 =  g0        & 0x0F0F0F0Fu;   qw.s1 = (g0 >> 4) & 0x0F0F0F0Fu;
+            qw.s2 =  g1        & 0x0F0F0F0Fu;   qw.s3 = (g1 >> 4) & 0x0F0F0F0Fu;
+            qw.s4 =  g2        & 0x0F0F0F0Fu;   qw.s5 = (g2 >> 4) & 0x0F0F0F0Fu;
+            qw.s6 =  g3        & 0x0F0F0F0Fu;   qw.s7 = (g3 >> 4) & 0x0F0F0F0Fu;
+        }
+
+        for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) {
+            const uint t = idx >> 1;
+            const uint v = idx & 1;
+            const uint c = col_base + t;
+            sh_qa4[t][v] = (c < (uint)n_no_padding)
+                         ? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2))
+                         : (uint4)(0u);
+        }
+        if (lid < TILESIZE_N) {
+            const uint c = col_base + lid;
+            const bool ok = c < (uint)n_no_padding;
+            sh_da[lid] = ok ? src1_da[c * k_b + sub] : (half)0;
+            sh_sa[lid] = ok ? src1_sa[c * k_b + sub] : (half)0;
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3]))
+        #pragma unroll
+        for (int g = 0; g < NGROUPS; ++g) {
+            const int b = g * 4;
+            float4 rf;
+#define TDOT4(T) dot8_q8a_v(qw, (uint4)(sh_qa4[T][0]), (uint4)(sh_qa4[T][1]))
+            rf.s0 = (float)TDOT4(b+0);  rf.s1 = (float)TDOT4(b+1);
+            rf.s2 = (float)TDOT4(b+2);  rf.s3 = (float)TDOT4(b+3);
+            acc[g] += dl * (LD4(sh_da, b) * rf) + (dl * dlt) * LD4(sh_sa, b);
+        }
+#undef TDOT4
+#undef LD4
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    if (!row_valid) {
+        return;
+    }
+
+    #pragma unroll
+    for (int g = 0; g < NGROUPS; ++g) {
+        const uint b = (uint)(g * 4);
+        const float4 a = acc[g];
+        const uint c0 = col_base + b;
+        if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0;
+        if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1;
+        if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2;
+        if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3;
+    }
+#undef NGROUPS
+#undef IQ1_S_GRID
+}

ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq1_m_f32.cl

@@ -0,0 +1,134 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define QK_K 256
+#define IQ1_M_DELTA 0.125f
+
+#ifndef IQ1_M_MV_NSG
+#define IQ1_M_MV_NSG 4
+#endif
+
+inline float4 iq1_m_vals(uint g) {
+    return (float4)((float)( g        & 0xFu), (float)((g >>  8) & 0xFu),
+                    (float)((g >> 16) & 0xFu), (float)((g >> 24) & 0xFu));
+}
+
+inline float iq1_m_super(uint w0, uint w1, uint w2, uint w3) {
+    const ushort u16 = (ushort)(((w0 >> 12)) | ((w1 >> 8) & 0x00f0u)
+                              | ((w2 >> 4) & 0x0f00u) | (w3 & 0xf000u));
+    return (float)as_half(u16);
+}
+
+kernel void kernel_gemv_noshuffle_iq1_m_f32(
+        __read_only image1d_buffer_t grid_img,
+        global const uchar  * src0_qs,
+        global const uchar  * src0_qh,
+        global const ushort * src0_sc,
+        global const float  * src1,
+        ulong offset1,
+        global float * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        int ne10,
+        int ne0
+) {
+    src1 = (global const float *)((global const char *)src1 + offset1);
+    dst  = (global float       *)((global char       *)dst  + offsetd);
+
+    const uint m   = (uint)ne01;
+    const uint K   = (uint)ne00;
+    const uint nsb = K / QK_K;
+
+    const uint lid = get_local_id(0);
+    const uint sgi = get_local_id(1);
+    const uint col = get_group_id(1);
+
+    global const float * y = src1 + (ulong)col * (uint)ne10;
+
+#define IQ1_M_GRID(i) (read_imageui(grid_img, (int)(i)).x)
+
+    const uint mh  = m >> 1;
+    const uint j   = get_group_id(0) * 64u + lid;
+    const uint row = j << 1;
+
+    float sumf = 0.f, sumf1 = 0.f;
+
+    if (j < mh) {
+        global const ushort * qsu = (global const ushort *)src0_qs;
+        global const ushort * qhu = (global const ushort *)src0_qh;
+        global const uint   * scu = (global const uint   *)src0_sc;
+
+        for (uint ib = sgi; ib < nsb; ib += IQ1_M_MV_NSG) {
+            const uint sb = j + (ib * 4u) * mh;
+            const uint w0 = scu[sb + 0u * mh];
+            const uint w1 = scu[sb + 1u * mh];
+            const uint w2 = scu[sb + 2u * mh];
+            const uint w3 = scu[sb + 3u * mh];
+            const float d0 = iq1_m_super( w0        & 0xFFFFu,  w1        & 0xFFFFu,
+                                         w2        & 0xFFFFu,  w3        & 0xFFFFu);
+            const float d1 = iq1_m_super((w0 >> 16) & 0xFFFFu, (w1 >> 16) & 0xFFFFu,
+                                        (w2 >> 16) & 0xFFFFu, (w3 >> 16) & 0xFFFFu);
+
+            float acc0 = 0.f, acc1 = 0.f;
+
+            for (uint sub_l = 0; sub_l < 8u; ++sub_l) {
+                const uint sub = ib * 8u + sub_l;
+                const uint scw = scu[j + (ib * 4u + (sub_l >> 1)) * mh];
+                const uint grp = sub * 8u;
+                const uint shb = 3u * (2u * (sub_l & 1u));
+
+                for (uint l = 0; l < 4u; ++l) {
+                    const uint sh  = shb + 3u * (l >> 1);
+                    const float s0 = (float)(2u * (((scw       ) >> sh) & 7u) + 1u);
+                    const float s1 = (float)(2u * (((scw >> 16) >> sh) & 7u) + 1u);
+
+                    const uint qhv = (uint)qhu[j + (sub * 2u + (l >> 1)) * mh];
+                    const uint h0  =  qhv       & 0xFFu;
+                    const uint h1  = (qhv >> 8) & 0xFFu;
+                    const uint qsv = (uint)qsu[j + (sub * 4u + l) * mh];
+
+                    const uint msk = 0x08u << (4u * (l & 1u));
+                    const float t0 = (h0 & msk) ? (-1.f - IQ1_M_DELTA) : (-1.f + IQ1_M_DELTA);
+                    const float t1 = (h1 & msk) ? (-1.f - IQ1_M_DELTA) : (-1.f + IQ1_M_DELTA);
+
+                    const uint g0 = IQ1_M_GRID(( qsv       & 0xFFu)
+                                              | ((((h0 >> (4u*(l & 1u))) & 7u) << 8)));
+                    const uint g1 = IQ1_M_GRID(((qsv >> 8) & 0xFFu)
+                                              | ((((h1 >> (4u*(l & 1u))) & 7u) << 8)));
+
+                    const float4 y0 = vload4(grp + 2u*l + 0u, y);
+                    const float4 y1 = vload4(grp + 2u*l + 1u, y);
+
+                    const float4 ys = y0 + y1;
+                    const float  as = ys.s0 + ys.s1 + ys.s2 + ys.s3;
+
+                    const float a0 = dot(y0, iq1_m_vals(g0)) + dot(y1, iq1_m_vals(g0 >> 4));
+                    const float a1 = dot(y0, iq1_m_vals(g1)) + dot(y1, iq1_m_vals(g1 >> 4));
+                    acc0 += s0 * (a0 + t0 * as);
+                    acc1 += s1 * (a1 + t1 * as);
+                }
+            }
+            sumf  += d0 * acc0;
+            sumf1 += d1 * acc1;
+        }
+    }
+
+#if IQ1_M_MV_NSG > 1
+    __local float2 part[IQ1_M_MV_NSG][64];
+    part[sgi][lid] = (float2)(sumf, sumf1);
+    barrier(CLK_LOCAL_MEM_FENCE);
+    if (sgi != 0) {
+        return;
+    }
+    for (uint s = 1; s < IQ1_M_MV_NSG; ++s) {
+        const float2 p = part[s][lid];
+        sumf  += p.s0;
+        sumf1 += p.s1;
+    }
+#endif
+
+    if (j < mh) {
+        vstore2((float2)(sumf, sumf1), 0, dst + (ulong)col * (uint)ne0 + row);
+    }
+#undef IQ1_M_GRID
+}

ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq1_s_f32.cl

@@ -0,0 +1,117 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define QK_K 256
+
+#ifndef IQ1_S_MV_NSG
+#define IQ1_S_MV_NSG 4
+#endif
+
+inline float4 iq1_s_vals(uint g) {
+    return (float4)((float)( g        & 0xFu), (float)((g >>  8) & 0xFu),
+                    (float)((g >> 16) & 0xFu), (float)((g >> 24) & 0xFu));
+}
+
+kernel void kernel_gemv_noshuffle_iq1_s_f32(
+        __read_only image1d_buffer_t grid_img,
+        global const uchar  * src0_qs,
+        global const ushort * src0_qh,
+        global const half   * src0_d,
+        global const float  * src1,
+        ulong offset1,
+        global float * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        int ne10,
+        int ne0
+) {
+    src1 = (global const float *)((global const char *)src1 + offset1);
+    dst  = (global float       *)((global char       *)dst  + offsetd);
+
+    const uint m   = (uint)ne01;
+    const uint K   = (uint)ne00;
+    const uint nsb = K / QK_K;
+
+    const uint lid = get_local_id(0);
+    const uint sgi = get_local_id(1);
+    const uint col = get_group_id(1);
+
+    global const float * y = src1 + (ulong)col * (uint)ne10;
+
+#define IQ1_S_GRID(i) (read_imageui(grid_img, (int)(i)).x)
+
+    const uint mh  = m >> 1;
+    const uint j   = get_group_id(0) * 64u + lid;
+    const uint row = j << 1;
+
+    float sumf = 0.f, sumf1 = 0.f;
+
+    if (j < mh) {
+        global const ushort * qsu = (global const ushort *)src0_qs;
+
+        for (uint ib = sgi; ib < nsb; ib += IQ1_S_MV_NSG) {
+            const half2 dh = vload2(j + ib * mh, src0_d);
+
+            float acc0 = 0.f, acc1 = 0.f;
+            for (uint sb = 0; sb < 8u; ++sb) {
+                const uint  sub = ib * 8u + sb;
+                const ushort2 qhv = vload2(j + sub * mh, src0_qh);
+                const uint  h0  = (uint)qhv.s0, h1 = (uint)qhv.s1;
+
+                const float s0  = (float)(2u * ((h0 >> 12) & 7u) + 1u);
+                const float s1  = (float)(2u * ((h1 >> 12) & 7u) + 1u);
+                const float t0  = ((h0 & 0x8000u) ? -0.125f : 0.125f) - 1.0f;
+                const float t1  = ((h1 & 0x8000u) ? -0.125f : 0.125f) - 1.0f;
+
+                const uint gb  = j + (sub * 4u) * mh;
+                const uint grp = sub * 8u;
+
+                float a0 = 0.f, a1 = 0.f;
+                float4 as = (float4)(0.f);
+                for (uint l = 0; l < 4u; ++l) {
+                    const uint qsv = (uint)qsu[gb + l * mh];
+                    const uint g0  = IQ1_S_GRID(( qsv       & 0xFFu) | ((((h0 >> (3u*l)) & 7u) << 8)));
+                    const uint g1  = IQ1_S_GRID(((qsv >> 8) & 0xFFu) | ((((h1 >> (3u*l)) & 7u) << 8)));
+
+                    const float4 y0 = vload4(grp + 2u*l + 0u, y);
+                    const float4 y1 = vload4(grp + 2u*l + 1u, y);
+                    as += y0 + y1;
+                    a0 += dot(y0, iq1_s_vals(g0)) + dot(y1, iq1_s_vals(g0 >> 4));
+                    a1 += dot(y0, iq1_s_vals(g1)) + dot(y1, iq1_s_vals(g1 >> 4));
+                }
+
+                const float asum = as.s0 + as.s1 + as.s2 + as.s3;
+                acc0 += s0 * (a0 + t0 * asum);
+                acc1 += s1 * (a1 + t1 * asum);
+            }
+            sumf  += (float)dh.s0 * acc0;
+            sumf1 += (float)dh.s1 * acc1;
+        }
+    }
+
+#if IQ1_S_MV_NSG > 1
+    __local float2 part[IQ1_S_MV_NSG][64];
+    part[sgi][lid] = (float2)(sumf, sumf1);
+    barrier(CLK_LOCAL_MEM_FENCE);
+    if (sgi != 0) {
+        return;
+    }
+    for (uint s = 1; s < IQ1_S_MV_NSG; ++s) {
+        const float2 p = part[s][lid];
+        sumf  += p.s0;
+        sumf1 += p.s1;
+    }
+#endif
+
+    if (j < mh) {
+        vstore2((float2)(sumf, sumf1), 0, dst + (ulong)col * (uint)ne0 + row);
+    }
+#undef IQ1_S_GRID
+}
+
+kernel void kernel_iq1_s_grid_export(global uint * out) {
+    const uint i = get_global_id(0);
+    if (i < 2048u) {
+        out[i] = iq1_s_grid[i];
+    }
+}

ggml/src/ggml-opencl/kernels/iq1_s_grid.cl

@@ -0,0 +1,259 @@
+// IQ1 codebook: each byte packs grid[j] + 1 and grid[j + 4] + 1.
+constant uint iq1_s_grid[2048] = {
+    0x00000000, 0x00000002, 0x00000101, 0x00000200, 0x00000202, 0x00010001, 0x00010101, 0x00020000,
+    0x00020002, 0x00020200, 0x00020202, 0x01000101, 0x01010001, 0x01010100, 0x01010102, 0x01020101,
+    0x02000000, 0x02000002, 0x02000200, 0x02000202, 0x02010101, 0x02020000, 0x02020002, 0x02020200,
+    0x02020202, 0x00000110, 0x00000111, 0x00010011, 0x00010110, 0x00010112, 0x00010211, 0x00010212,
+    0x00020111, 0x01000011, 0x01000112, 0x01000211, 0x01010012, 0x01010111, 0x01010212, 0x01020011,
+    0x01020110, 0x01020112, 0x01020210, 0x02000111, 0x02010011, 0x02010110, 0x02010112, 0x02020111,
+    0x00000020, 0x00000022, 0x00000220, 0x00000222, 0x00010121, 0x00020020, 0x00020022, 0x00020220,
+    0x00020222, 0x01000121, 0x01010021, 0x01010221, 0x01020120, 0x01020221, 0x02000020, 0x02000022,
+    0x02000220, 0x02000222, 0x02010021, 0x02010121, 0x02010221, 0x02020020, 0x02020022, 0x02020220,
+    0x02020222, 0x00011001, 0x00011100, 0x00011102, 0x00021101, 0x01001001, 0x01001201, 0x01011101,
+    0x01011202, 0x01021100, 0x01021101, 0x02011001, 0x02011201, 0x02021101, 0x00001011, 0x00001110,
+    0x00001111, 0x00001112, 0x00011111, 0x00011210, 0x00011212, 0x00021211, 0x01001010, 0x01001111,
+    0x01001212, 0x01011010, 0x01011011, 0x01011110, 0x01011111, 0x01011112, 0x01011211, 0x01021010,
+    0x01021012, 0x01021111, 0x01021210, 0x01021212, 0x02001011, 0x02011011, 0x02011111, 0x02011210,
+    0x02011212, 0x02021011, 0x02021110, 0x02021111, 0x02021112, 0x02021211, 0x00011120, 0x00011221,
+    0x01001021, 0x01001120, 0x01011020, 0x01011022, 0x01011121, 0x01011220, 0x01021020, 0x01021021,
+    0x01021122, 0x01021221, 0x02001121, 0x02011021, 0x02011120, 0x02011221, 0x00002000, 0x00002002,
+    0x00002200, 0x00002202, 0x00012101, 0x00022000, 0x00022002, 0x00022200, 0x00022202, 0x01002101,
+    0x01012001, 0x01012102, 0x01022101, 0x02002000, 0x02002002, 0x02002200, 0x02002202, 0x02012101,
+    0x02022000, 0x02022002, 0x02022200, 0x02022202, 0x00002111, 0x00012011, 0x00012110, 0x00012211,
+    0x00022110, 0x00022111, 0x01002011, 0x01012010, 0x01012011, 0x01012111, 0x01022011, 0x01022110,
+    0x01022211, 0x02012011, 0x02012110, 0x02012112, 0x02012211, 0x02022111, 0x00002020, 0x00002022,
+    0x00002220, 0x00002222, 0x00012121, 0x00022020, 0x00022022, 0x00022220, 0x00022222, 0x01002121,
+    0x01012021, 0x01012221, 0x01022021, 0x01022121, 0x02002020, 0x02002022, 0x02002121, 0x02002220,
+    0x02002222, 0x02012121, 0x02022020, 0x02022022, 0x02022220, 0x02022222, 0x00110000, 0x00110001,
+    0x00110100, 0x00110201, 0x00120100, 0x00120101, 0x01100001, 0x01100100, 0x01110000, 0x01110101,
+    0x01110200, 0x01120001, 0x01120100, 0x01120101, 0x01120201, 0x02110001, 0x02110100, 0x02110102,
+    0x02120001, 0x02120101, 0x00100011, 0x00100110, 0x00100112, 0x00100211, 0x00110010, 0x00110012,
+    0x00110111, 0x00110210, 0x00120011, 0x00120110, 0x00120211, 0x01100111, 0x01100212, 0x01110010,
+    0x01110011, 0x01110012, 0x01110110, 0x01110111, 0x01110112, 0x01110211, 0x01120010, 0x01120111,
+    0x02100110, 0x02110012, 0x02110111, 0x02120011, 0x02120110, 0x00110021, 0x00110120, 0x00110122,
+    0x00120121, 0x01100020, 0x01100122, 0x01100221, 0x01110022, 0x01110121, 0x01110220, 0x01110222,
+    0x01120120, 0x01120122, 0x02100121, 0x02110021, 0x02110120, 0x02110122, 0x02120121, 0x00101001,
+    0x00101102, 0x00101201, 0x00111100, 0x00111101, 0x00111200, 0x00111201, 0x00121001, 0x00121102,
+    0x01101001, 0x01101101, 0x01101102, 0x01101200, 0x01101202, 0x01111001, 0x01111100, 0x01111101,
+    0x01111102, 0x01111201, 0x01121002, 0x01121101, 0x01121200, 0x02101100, 0x02101201, 0x02111000,
+    0x02111100, 0x02111101, 0x02111200, 0x02111201, 0x02111202, 0x02121001, 0x02121100, 0x02121101,
+    0x02121201, 0x00101012, 0x00101111, 0x00101212, 0x00111011, 0x00111110, 0x00111111, 0x00111112,
+    0x00111211, 0x00121010, 0x00121012, 0x00121111, 0x00121210, 0x00121212, 0x01101011, 0x01101110,
+    0x01101111, 0x01101112, 0x01111011, 0x01111012, 0x01111110, 0x01111111, 0x01111112, 0x01111211,
+    0x01111212, 0x01121011, 0x01121110, 0x01121111, 0x01121112, 0x01121211, 0x02101010, 0x02101012,
+    0x02101110, 0x02101111, 0x02101210, 0x02101212, 0x02111010, 0x02111011, 0x02111110, 0x02111111,
+    0x02111112, 0x02111211, 0x02111212, 0x02121010, 0x02121012, 0x02121111, 0x00101021, 0x00101120,
+    0x00101121, 0x00101122, 0x00111121, 0x00111122, 0x00111220, 0x00111222, 0x00121021, 0x00121122,
+    0x01101020, 0x01101022, 0x01101120, 0x01101121, 0x01101220, 0x01101222, 0x01111021, 0x01111121,
+    0x01111122, 0x01111220, 0x01111221, 0x01121021, 0x01121120, 0x01121121, 0x01121220, 0x01121221,
+    0x01121222, 0x02101122, 0x02101222, 0x02111022, 0x02111121, 0x02121120, 0x02121221, 0x00112001,
+    0x00112102, 0x00122101, 0x01102001, 0x01102100, 0x01102102, 0x01102201, 0x01112000, 0x01112101,
+    0x01112200, 0x01112202, 0x01122000, 0x01122001, 0x01122100, 0x01122102, 0x01122201, 0x02102101,
+    0x02112001, 0x02112100, 0x02122101, 0x00112010, 0x00112012, 0x00112111, 0x00112212, 0x00122011,
+    0x00122111, 0x01102012, 0x01102110, 0x01102111, 0x01102210, 0x01112011, 0x01112110, 0x01112111,
+    0x01112112, 0x01112211, 0x01112212, 0x01122010, 0x01122111, 0x01122212, 0x02102211, 0x02112011,
+    0x02112012, 0x02112111, 0x02112210, 0x02122011, 0x02122112, 0x02122211, 0x00102221, 0x00112122,
+    0x00122120, 0x00122122, 0x01102120, 0x01102122, 0x01102221, 0x01112020, 0x01112022, 0x01112121,
+    0x01112220, 0x01122021, 0x01122122, 0x01122221, 0x02102121, 0x02112021, 0x02112122, 0x02112222,
+    0x00200000, 0x00200002, 0x00200200, 0x00200202, 0x00210101, 0x00220000, 0x00220002, 0x00220101,
+    0x00220200, 0x00220202, 0x01200101, 0x01210001, 0x01210201, 0x01220001, 0x01220101, 0x02200000,
+    0x02200002, 0x02200200, 0x02200202, 0x02210101, 0x02220000, 0x02220002, 0x02220101, 0x02220200,
+    0x02220202, 0x00200111, 0x00210011, 0x00210110, 0x00210211, 0x00220111, 0x01200012, 0x01200110,
+    0x01200211, 0x01210111, 0x01210210, 0x01210212, 0x01220011, 0x01220110, 0x01220111, 0x01220112,
+    0x02200111, 0x02210010, 0x02210112, 0x02210211, 0x02220111, 0x00200021, 0x00200220, 0x00200222,
+    0x00210021, 0x00210121, 0x00220020, 0x00220022, 0x00220220, 0x00220222, 0x01200121, 0x01210021,
+    0x01210122, 0x01210221, 0x01220121, 0x02200021, 0x02200220, 0x02200222, 0x02210021, 0x02210121,
+    0x02220020, 0x02220022, 0x02220220, 0x02220222, 0x00201101, 0x00211100, 0x00211102, 0x00211201,
+    0x00221101, 0x01201100, 0x01201101, 0x01201102, 0x01201201, 0x01211002, 0x01211101, 0x01211200,
+    0x01211202, 0x01221102, 0x02201101, 0x02211001, 0x02211100, 0x02211201, 0x02221001, 0x02221101,
+    0x00201211, 0x00211111, 0x00221011, 0x00221211, 0x01201010, 0x01201111, 0x01201210, 0x01211011,
+    0x01211110, 0x01211111, 0x01211211, 0x01221012, 0x01221111, 0x01221210, 0x02201211, 0x02211010,
+    0x02211110, 0x02211111, 0x02211210, 0x02211212, 0x02221011, 0x02221110, 0x02221112, 0x02221211,
+    0x00201121, 0x00211020, 0x00211022, 0x00211221, 0x00221121, 0x01201021, 0x01201221, 0x01211121,
+    0x01221020, 0x01221021, 0x01221221, 0x02201120, 0x02201122, 0x02211020, 0x02211222, 0x00202000,
+    0x00202002, 0x00202200, 0x00202202, 0x00212101, 0x00222000, 0x00222002, 0x00222200, 0x00222202,
+    0x01202101, 0x01212001, 0x01212100, 0x01222101, 0x02202000, 0x02202002, 0x02202200, 0x02202202,
+    0x02222000, 0x02222002, 0x02222200, 0x02222202, 0x00202211, 0x00212011, 0x00212110, 0x00212211,
+    0x00222111, 0x01202112, 0x01202211, 0x01212012, 0x01212111, 0x01222011, 0x01222110, 0x01222112,
+    0x01222211, 0x02202111, 0x02212010, 0x02212112, 0x02212211, 0x02222110, 0x02222111, 0x00202020,
+    0x00202022, 0x00202220, 0x00202222, 0x00222020, 0x00222022, 0x00222220, 0x00222222, 0x01202121,
+    0x01212021, 0x01212122, 0x01212221, 0x01222121, 0x02202020, 0x02202022, 0x02202220, 0x02202222,
+    0x02212121, 0x02222020, 0x02222022, 0x02222220, 0x02222222, 0x10000101, 0x10010001, 0x10010102,
+    0x10020101, 0x11000201, 0x11010002, 0x11010101, 0x11010200, 0x11010202, 0x11020001, 0x11020100,
+    0x11020102, 0x12010100, 0x12010201, 0x12020001, 0x12020102, 0x10000010, 0x10000011, 0x10000110,
+    0x10000112, 0x10000211, 0x10010012, 0x10010111, 0x10010112, 0x10010210, 0x10010212, 0x10020011,
+    0x10020112, 0x10020211, 0x11000111, 0x11000210, 0x11000212, 0x11010011, 0x11010110, 0x11010111,
+    0x11010112, 0x11010211, 0x11010212, 0x11020111, 0x11020210, 0x11020212, 0x12000011, 0x12000110,
+    0x12000112, 0x12010010, 0x12010012, 0x12010111, 0x12020010, 0x12020011, 0x12020012, 0x10000121,
+    0x10010021, 0x10010120, 0x10010122, 0x10020121, 0x11000021, 0x11010022, 0x11010121, 0x11010222,
+    0x11020120, 0x11020221, 0x12000221, 0x12010120, 0x12020121, 0x10001001, 0x10011101, 0x10011201,
+    0x10021201, 0x11001101, 0x11001200, 0x11001202, 0x11011001, 0x11011100, 0x11011101, 0x11011102,
+    0x11021001, 0x11021002, 0x11021101, 0x11021200, 0x11021202, 0x12001001, 0x12001102, 0x12001201,
+    0x12011000, 0x12011002, 0x12011101, 0x12021000, 0x12021001, 0x12021201, 0x10001011, 0x10001012,
+    0x10001111, 0x10001212, 0x10011011, 0x10011110, 0x10011111, 0x10011112, 0x10011211, 0x10021010,
+    0x10021111, 0x10021212, 0x11001011, 0x11001110, 0x11001111, 0x11001112, 0x11001211, 0x11011010,
+    0x11011011, 0x11011110, 0x11011111, 0x11011112, 0x11011210, 0x11011211, 0x11021011, 0x11021110,
+    0x11021111, 0x11021112, 0x11021211, 0x12001012, 0x12001110, 0x12001111, 0x12001210, 0x12011011,
+    0x12011110, 0x12011111, 0x12011112, 0x12011211, 0x12011212, 0x12021111, 0x12021210, 0x12021212,
+    0x10001021, 0x10001121, 0x10001221, 0x10011120, 0x10011121, 0x10011220, 0x10011222, 0x10021021,
+    0x10021120, 0x10021221, 0x11001020, 0x11001022, 0x11001121, 0x11001220, 0x11011020, 0x11011021,
+    0x11011022, 0x11011121, 0x11011122, 0x11011221, 0x11021022, 0x11021121, 0x11021220, 0x12001021,
+    0x12001121, 0x12001222, 0x12011120, 0x12011121, 0x12021021, 0x12021120, 0x12021122, 0x10002101,
+    0x10012001, 0x10012101, 0x10012202, 0x10022101, 0x11002002, 0x11002201, 0x11012000, 0x11012101,
+    0x11012200, 0x11022001, 0x11022100, 0x11022102, 0x11022201, 0x12002101, 0x12012001, 0x12012100,
+    0x12012102, 0x12012201, 0x12022101, 0x10002011, 0x10002111, 0x10002112, 0x10002212, 0x10012010,
+    0x10012110, 0x10012111, 0x10012210, 0x10022011, 0x10022110, 0x10022112, 0x11002010, 0x11002111,
+    0x11002212, 0x11012011, 0x11012012, 0x11012110, 0x11012111, 0x11012112, 0x11012211, 0x11022010,
+    0x11022012, 0x11022111, 0x11022112, 0x11022212, 0x12002112, 0x12002211, 0x12012012, 0x12012111,
+    0x12012112, 0x12012210, 0x12022011, 0x12022110, 0x12022112, 0x12022211, 0x10012122, 0x11002120,
+    0x11002122, 0x11002221, 0x11012121, 0x11012220, 0x11012222, 0x11022120, 0x11022221, 0x12012120,
+    0x12022121, 0x10100001, 0x10100100, 0x10100101, 0x10100102, 0x10100201, 0x10110002, 0x10110101,
+    0x10110202, 0x10120001, 0x10120100, 0x10120201, 0x11100000, 0x11100101, 0x11100200, 0x11110001,
+    0x11110100, 0x11110101, 0x11110102, 0x11110201, 0x11120101, 0x11120200, 0x12100102, 0x12100201,
+    0x12110101, 0x12110200, 0x12120000, 0x12120001, 0x12120102, 0x12120201, 0x10100111, 0x10100210,
+    0x10100211, 0x10100212, 0x10110011, 0x10110110, 0x10110111, 0x10110112, 0x10110210, 0x10110211,
+    0x10120010, 0x10120111, 0x10120112, 0x10120210, 0x10120212, 0x11100011, 0x11100110, 0x11100111,
+    0x11100112, 0x11100211, 0x11110010, 0x11110011, 0x11110012, 0x11110110, 0x11110111, 0x11110112,
+    0x11110210, 0x11110211, 0x11110212, 0x11120011, 0x11120110, 0x11120111, 0x11120112, 0x11120211,
+    0x12100012, 0x12100111, 0x12110011, 0x12110110, 0x12110111, 0x12110112, 0x12110211, 0x12120010,
+    0x12120111, 0x12120212, 0x10100021, 0x10100122, 0x10110022, 0x10110121, 0x10110222, 0x10120021,
+    0x10120120, 0x11100022, 0x11100121, 0x11100222, 0x11110021, 0x11110120, 0x11110121, 0x11110122,
+    0x11110221, 0x11120022, 0x11120121, 0x12100121, 0x12110020, 0x12110022, 0x12110121, 0x12110221,
+    0x12110222, 0x12120120, 0x10101100, 0x10101101, 0x10111001, 0x10111100, 0x10111101, 0x10111102,
+    0x10111200, 0x10111201, 0x10121001, 0x10121101, 0x10121200, 0x10121202, 0x11101001, 0x11101100,
+    0x11101101, 0x11101102, 0x11101201, 0x11101202, 0x11111000, 0x11111001, 0x11111100, 0x11111101,
+    0x11111102, 0x11111200, 0x11111201, 0x11111202, 0x11121001, 0x11121002, 0x11121100, 0x11121101,
+    0x11121102, 0x11121201, 0x12101000, 0x12101200, 0x12101202, 0x12111001, 0x12111100, 0x12111101,
+    0x12111102, 0x12111201, 0x12121001, 0x12121100, 0x12121101, 0x12121202, 0x10101011, 0x10101012,
+    0x10101110, 0x10101111, 0x10101112, 0x10101211, 0x10111010, 0x10111011, 0x10111012, 0x10111110,
+    0x10111111, 0x10111112, 0x10111211, 0x10111212, 0x10121011, 0x10121110, 0x10121111, 0x10121112,
+    0x10121211, 0x11101010, 0x11101011, 0x11101012, 0x11101110, 0x11101111, 0x11101112, 0x11101210,
+    0x11101211, 0x11111010, 0x11111011, 0x11111012, 0x11111110, 0x11111111, 0x11111112, 0x11111210,
+    0x11111211, 0x11111212, 0x11121010, 0x11121011, 0x11121110, 0x11121111, 0x11121112, 0x11121210,
+    0x11121211, 0x11121212, 0x12101011, 0x12101110, 0x12101111, 0x12101211, 0x12101212, 0x12111010,
+    0x12111011, 0x12111110, 0x12111111, 0x12111112, 0x12111210, 0x12111211, 0x12121011, 0x12121110,
+    0x12121111, 0x12121112, 0x12121211, 0x10101020, 0x10101021, 0x10101022, 0x10101120, 0x10101122,
+    0x10101220, 0x10101221, 0x10111021, 0x10111120, 0x10111121, 0x10111220, 0x10111221, 0x10121020,
+    0x10121021, 0x10121022, 0x10121120, 0x10121121, 0x10121122, 0x10121220, 0x10121221, 0x11101021,
+    0x11101121, 0x11101122, 0x11101220, 0x11101221, 0x11101222, 0x11111020, 0x11111021, 0x11111022,
+    0x11111120, 0x11111121, 0x11111122, 0x11111220, 0x11111221, 0x11111222, 0x11121021, 0x11121120,
+    0x11121121, 0x11121221, 0x12101022, 0x12101121, 0x12101122, 0x12101220, 0x12101221, 0x12101222,
+    0x12111021, 0x12111121, 0x12111222, 0x12121022, 0x12121121, 0x12121122, 0x12121220, 0x12121221,
+    0x10102100, 0x10102101, 0x10102102, 0x10102201, 0x10112000, 0x10112101, 0x10112200, 0x10122001,
+    0x10122202, 0x11102101, 0x11102200, 0x11102202, 0x11112001, 0x11112100, 0x11112101, 0x11112102,
+    0x11112200, 0x11112201, 0x11122000, 0x11122002, 0x11122100, 0x11122101, 0x12102002, 0x12102201,
+    0x12112000, 0x12112002, 0x12112101, 0x12112200, 0x12122001, 0x12122201, 0x10102011, 0x10102012,
+    0x10102111, 0x10102212, 0x10112011, 0x10112110, 0x10112111, 0x10112112, 0x10112211, 0x10122111,
+    0x11102011, 0x11102110, 0x11102111, 0x11102112, 0x11102211, 0x11112010, 0x11112011, 0x11112012,
+    0x11112110, 0x11112111, 0x11112112, 0x11112210, 0x11112211, 0x11112212, 0x11122011, 0x11122110,
+    0x11122111, 0x11122112, 0x11122211, 0x12102011, 0x12102111, 0x12102211, 0x12112011, 0x12112110,
+    0x12112111, 0x12112112, 0x12112210, 0x12112211, 0x12122111, 0x10102120, 0x10102220, 0x10112121,
+    0x10112222, 0x10122020, 0x10122121, 0x10122122, 0x10122221, 0x11102121, 0x11102220, 0x11102221,
+    0x11112021, 0x11112121, 0x11112122, 0x11112220, 0x11112221, 0x11122022, 0x11122121, 0x11122220,
+    0x11122222, 0x12102021, 0x12102222, 0x12112022, 0x12112121, 0x12112122, 0x12112220, 0x12112222,
+    0x12122021, 0x10200101, 0x10210100, 0x10210102, 0x10210201, 0x10220101, 0x11200100, 0x11210000,
+    0x11210101, 0x11210102, 0x11210200, 0x11210202, 0x11220001, 0x11220100, 0x11220102, 0x11220201,
+    0x12200001, 0x12210102, 0x12220101, 0x10200011, 0x10200110, 0x10200112, 0x10200211, 0x10210012,
+    0x10210111, 0x10220011, 0x10220012, 0x10220112, 0x10220211, 0x11200111, 0x11200211, 0x11210011,
+    0x11210111, 0x11210112, 0x11210211, 0x11220111, 0x11220112, 0x11220212, 0x12200110, 0x12200212,
+    0x12210012, 0x12210111, 0x12220011, 0x12220112, 0x12220211, 0x10210021, 0x10210122, 0x10210221,
+    0x11200020, 0x11200021, 0x11200122, 0x11210121, 0x11210122, 0x11210220, 0x11220020, 0x12200121,
+    0x12210021, 0x12210122, 0x12220121, 0x10211001, 0x10211002, 0x10211101, 0x10211102, 0x10211202,
+    0x10221001, 0x10221102, 0x10221201, 0x11201000, 0x11201002, 0x11201101, 0x11201200, 0x11201202,
+    0x11211001, 0x11211100, 0x11211101, 0x11211102, 0x11211201, 0x11211202, 0x11221000, 0x11221002,
+    0x11221101, 0x12201100, 0x12201101, 0x12201201, 0x12211000, 0x12211002, 0x12211100, 0x12211101,
+    0x12211102, 0x12211200, 0x12211202, 0x12221001, 0x12221100, 0x12221201, 0x10201111, 0x10201210,
+    0x10201212, 0x10211011, 0x10211111, 0x10211112, 0x10211211, 0x11201110, 0x11201111, 0x11201112,
+    0x11201211, 0x11211010, 0x11211011, 0x11211110, 0x11211111, 0x11211112, 0x11211211, 0x11221011,
+    0x11221110, 0x11221111, 0x11221112, 0x11221211, 0x12201112, 0x12201211, 0x12201212, 0x12211011,
+    0x12211111, 0x12211112, 0x12211211, 0x12211212, 0x12221012, 0x12221111, 0x12221112, 0x12221210,
+    0x10201022, 0x10201221, 0x10211121, 0x10221020, 0x10221122, 0x10221220, 0x10221221, 0x11201020,
+    0x11201121, 0x11201220, 0x11201222, 0x11211021, 0x11211120, 0x11211121, 0x11211122, 0x11211220,
+    0x11211222, 0x11221020, 0x11221121, 0x11221220, 0x12201020, 0x12201022, 0x12201121, 0x12201222,
+    0x12211120, 0x12211122, 0x12211220, 0x12211221, 0x12221020, 0x12221120, 0x12221122, 0x12221222,
+    0x10212102, 0x10212201, 0x10222101, 0x11202001, 0x11212002, 0x11212101, 0x11212202, 0x11222001,
+    0x11222201, 0x12202101, 0x12212001, 0x12212200, 0x12222102, 0x10202011, 0x10202110, 0x10212010,
+    0x10212111, 0x10222011, 0x10222110, 0x10222112, 0x10222211, 0x11202010, 0x11202011, 0x11202111,
+    0x11202112, 0x11202210, 0x11212011, 0x11212110, 0x11212111, 0x11212112, 0x11212211, 0x11222010,
+    0x11222111, 0x11222212, 0x12202012, 0x12202110, 0x12202212, 0x12212111, 0x12222011, 0x12222110,
+    0x12222111, 0x12222211, 0x10212021, 0x10212122, 0x10212220, 0x11202021, 0x11202120, 0x11202221,
+    0x11212020, 0x11212121, 0x11212220, 0x11212222, 0x11222120, 0x11222121, 0x11222221, 0x12202122,
+    0x12212120, 0x12212220, 0x12212222, 0x12222122, 0x20000000, 0x20000002, 0x20000200, 0x20000202,
+    0x20020000, 0x20020002, 0x20020200, 0x20020202, 0x21000101, 0x21010000, 0x21010001, 0x21010100,
+    0x21010102, 0x21010201, 0x21020101, 0x22000000, 0x22000002, 0x22000200, 0x22000202, 0x22010101,
+    0x22020000, 0x22020002, 0x22020200, 0x22020202, 0x20000111, 0x20010011, 0x20010110, 0x20010112,
+    0x20010211, 0x20020111, 0x21000011, 0x21000110, 0x21000211, 0x21010010, 0x21010012, 0x21010111,
+    0x21010112, 0x21010210, 0x21010211, 0x21020110, 0x21020112, 0x21020211, 0x22000111, 0x22000211,
+    0x22010110, 0x22010112, 0x22010211, 0x22020111, 0x20000020, 0x20000022, 0x20000220, 0x20000222,
+    0x20010121, 0x20020020, 0x20020022, 0x20020220, 0x20020222, 0x21010021, 0x21010120, 0x21010221,
+    0x21020121, 0x22000020, 0x22000022, 0x22000220, 0x22000222, 0x22010121, 0x22020020, 0x22020022,
+    0x22020220, 0x22020222, 0x20011100, 0x20011201, 0x21001001, 0x21001100, 0x21011001, 0x21011101,
+    0x21011202, 0x21021001, 0x21021100, 0x21021201, 0x22011100, 0x22011201, 0x20001011, 0x20001211,
+    0x20011012, 0x20011111, 0x20011212, 0x20021112, 0x20021211, 0x21001010, 0x21001011, 0x21001111,
+    0x21001210, 0x21011011, 0x21011110, 0x21011111, 0x21011112, 0x21011211, 0x21011212, 0x21021111,
+    0x21021112, 0x21021210, 0x21021212, 0x22001011, 0x22001110, 0x22001112, 0x22001211, 0x22011010,
+    0x22011012, 0x22011111, 0x22011210, 0x22021112, 0x20011021, 0x20011122, 0x20011221, 0x20021121,
+    0x21001021, 0x21001120, 0x21001221, 0x21001222, 0x21011020, 0x21011121, 0x21011221, 0x21011222,
+    0x21021021, 0x21021122, 0x21021222, 0x22001121, 0x22011021, 0x22011222, 0x22021120, 0x20002000,
+    0x20002002, 0x20002200, 0x20002202, 0x20012101, 0x20022000, 0x20022002, 0x20022200, 0x20022202,
+    0x21002001, 0x21002101, 0x21012001, 0x21012100, 0x21012201, 0x21022101, 0x21022201, 0x22002000,
+    0x22002002, 0x22002200, 0x22002202, 0x22012101, 0x22022000, 0x22022002, 0x22022200, 0x22022202,
+    0x20002111, 0x20002112, 0x20012011, 0x20012110, 0x20012112, 0x20022111, 0x21002011, 0x21002110,
+    0x21002112, 0x21002211, 0x21012010, 0x21012012, 0x21012111, 0x21012212, 0x21022011, 0x21022110,
+    0x22002111, 0x22012112, 0x22012211, 0x22022111, 0x20002020, 0x20002022, 0x20002220, 0x20002222,
+    0x20012121, 0x20022020, 0x20022022, 0x20022220, 0x20022222, 0x21002121, 0x21012021, 0x21012120,
+    0x21012122, 0x22002020, 0x22002022, 0x22002220, 0x22002222, 0x22012121, 0x22022020, 0x22022022,
+    0x22022220, 0x22022222, 0x20100101, 0x20110001, 0x20110102, 0x20110200, 0x20110201, 0x20120101,
+    0x21100001, 0x21100102, 0x21100201, 0x21110101, 0x21110200, 0x21110202, 0x21120201, 0x21120202,
+    0x22100101, 0x22110001, 0x22110100, 0x22110102, 0x22110201, 0x22120101, 0x20100011, 0x20100110,
+    0x20100112, 0x20100211, 0x20110010, 0x20110111, 0x20110210, 0x20110212, 0x20120011, 0x20120110,
+    0x20120112, 0x20120211, 0x21100010, 0x21100111, 0x21110010, 0x21110011, 0x21110110, 0x21110111,
+    0x21110112, 0x21110211, 0x21120012, 0x21120111, 0x22100110, 0x22100112, 0x22110012, 0x22110111,
+    0x22110210, 0x22120011, 0x22120110, 0x22120112, 0x22120211, 0x20100121, 0x20110021, 0x20110120,
+    0x20110221, 0x20120121, 0x21100120, 0x21100122, 0x21100221, 0x21110020, 0x21110022, 0x21110121,
+    0x21110220, 0x21120122, 0x21120221, 0x22100121, 0x22110120, 0x22110122, 0x22120221, 0x20101001,
+    0x20101100, 0x20101102, 0x20111000, 0x20111101, 0x20111200, 0x20121102, 0x21101000, 0x21101202,
+    0x21111001, 0x21111100, 0x21111101, 0x21111102, 0x21111200, 0x21111201, 0x21121000, 0x21121001,
+    0x21121002, 0x21121101, 0x22101100, 0x22101102, 0x22111002, 0x22111100, 0x22111101, 0x22111200,
+    0x22121001, 0x22121201, 0x20101010, 0x20101111, 0x20101210, 0x20101212, 0x20111010, 0x20111011,
+    0x20111110, 0x20111111, 0x20111112, 0x20111211, 0x20121011, 0x20121111, 0x20121211, 0x20121212,
+    0x21101011, 0x21101110, 0x21101111, 0x21101112, 0x21101211, 0x21111010, 0x21111011, 0x21111012,
+    0x21111110, 0x21111111, 0x21111112, 0x21111210, 0x21111211, 0x21111212, 0x21121011, 0x21121110,
+    0x21121111, 0x21121112, 0x21121211, 0x22101011, 0x22101111, 0x22101210, 0x22111011, 0x22111012,
+    0x22111110, 0x22111111, 0x22111112, 0x22111211, 0x22111212, 0x22121010, 0x22121012, 0x22121111,
+    0x22121210, 0x22121212, 0x20101021, 0x20101120, 0x20111020, 0x20111121, 0x20111221, 0x20121020,
+    0x20121122, 0x20121221, 0x21101121, 0x21101220, 0x21101221, 0x21111021, 0x21111022, 0x21111121,
+    0x21111122, 0x21111221, 0x21121121, 0x21121220, 0x22101022, 0x22101120, 0x22101221, 0x22101222,
+    0x22111022, 0x22111120, 0x22111121, 0x22121120, 0x22121122, 0x22121221, 0x20102101, 0x20112102,
+    0x20112201, 0x20122101, 0x21102001, 0x21102102, 0x21112000, 0x21112002, 0x21112101, 0x21112102,
+    0x21112202, 0x21122100, 0x21122101, 0x22102101, 0x22112001, 0x22112102, 0x22112201, 0x22122101,
+    0x20102110, 0x20102112, 0x20102211, 0x20112010, 0x20112012, 0x20112111, 0x20112210, 0x20112212,
+    0x20122010, 0x20122011, 0x20122110, 0x20122112, 0x21102010, 0x21102012, 0x21102111, 0x21102210,
+    0x21102212, 0x21112011, 0x21112110, 0x21112111, 0x21112112, 0x21112211, 0x21122012, 0x21122111,
+    0x21122112, 0x21122212, 0x22102011, 0x22102110, 0x22112010, 0x22112012, 0x22112111, 0x22112212,
+    0x22122011, 0x22122112, 0x20102121, 0x20112121, 0x20122121, 0x21102120, 0x21102122, 0x21102221,
+    0x21112020, 0x21112121, 0x21112220, 0x21122021, 0x22102121, 0x22112021, 0x22112120, 0x22112121,
+    0x22112122, 0x20200000, 0x20200002, 0x20200200, 0x20200202, 0x20210101, 0x20220000, 0x20220002,
+    0x20220200, 0x20220202, 0x21200101, 0x21210001, 0x21210100, 0x21210102, 0x21210201, 0x22200000,
+    0x22200002, 0x22200200, 0x22200202, 0x22210101, 0x22220000, 0x22220002, 0x22220200, 0x22220202,
+    0x20200111, 0x20200211, 0x20210011, 0x20210110, 0x20210112, 0x20210211, 0x20210212, 0x21200112,
+    0x21200211, 0x21210011, 0x21210111, 0x21210210, 0x21210212, 0x21220011, 0x21220110, 0x22200111,
+    0x22210010, 0x22210012, 0x22210112, 0x22210211, 0x20200022, 0x20200220, 0x20200222, 0x20210020,
+    0x20210221, 0x20220022, 0x20220220, 0x20220222, 0x21200121, 0x21210021, 0x21210122, 0x21210221,
+    0x21220121, 0x22200020, 0x22200022, 0x22200220, 0x22200222, 0x22210121, 0x22220020, 0x22220022,
+    0x22220220, 0x22220222, 0x20211201, 0x20221101, 0x21201001, 0x21201100, 0x21211000, 0x21211100,
+    0x21211101, 0x21211200, 0x21211202, 0x21221001, 0x21221101, 0x21221102, 0x21221200, 0x21221201,
+    0x22201101, 0x20201112, 0x20201211, 0x20211010, 0x20211012, 0x20211111, 0x20211210, 0x20221112,
+    0x20221211, 0x21201012, 0x21201111, 0x21211011, 0x21211110, 0x21211111, 0x21211112, 0x21211211,
+    0x21221111, 0x21221212, 0x22201011, 0x22201110, 0x22201111, 0x22201112, 0x22201211, 0x22211012,
+    0x22211111, 0x22211210, 0x20201121, 0x20211021, 0x20211122, 0x20211222, 0x20221021, 0x20221121,
+    0x21201120, 0x21201122, 0x21201222, 0x21211022, 0x21211121, 0x21211122, 0x21211220, 0x21221020,
+    0x21221022, 0x22201122, 0x22211020, 0x22211121, 0x22211122, 0x22211221, 0x22221021, 0x22221120,
+    0x22221122, 0x20202000, 0x20202002, 0x20202200, 0x20202202, 0x20222000, 0x20222002, 0x20222200,
+    0x20222202, 0x21212001, 0x21212100, 0x21212102, 0x21212201, 0x22202000, 0x22202002, 0x22202200,
+    0x22202202, 0x22212101, 0x22222000, 0x22222002, 0x22222200, 0x22222202, 0x20202111, 0x20212110,
+    0x20212211, 0x20222011, 0x20222111, 0x21202011, 0x21212010, 0x21212111, 0x21212212, 0x21222011,
+    0x21222112, 0x21222211, 0x22212010, 0x22212112, 0x20202020, 0x20202022, 0x20202220, 0x20202222,
+    0x20222020, 0x20222022, 0x20222220, 0x20222222, 0x21212021, 0x21212120, 0x21212122, 0x22202020,
+    0x22202022, 0x22202220, 0x22202222, 0x22212121, 0x22222020, 0x22222022, 0x22222220, 0x22222222
+};

ggml/src/ggml-opencl/kernels/mul_mm_iq1_m_f32_l4_lm.cl

@@ -0,0 +1,191 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_required_subgroup_size
+#define INTEL_GPU 1
+#endif
+
+#define QK_K 256
+
+struct block_iq1_m {
+    uchar qs[QK_K/8];
+    uchar qh[QK_K/16];
+    uchar scales[QK_K/32];
+};
+
+#define LOAD_VEC_A 4
+#define LOAD_VEC_B 4
+
+#define BM 64
+#define BN 64
+#define BK 32
+#ifdef INTEL_GPU
+#define TM 8
+#define TN 8
+#else
+#define TM 4
+#define TN 8
+#endif
+
+kernel void kernel_mul_mm_iq1_m_f32_l4_lm(
+    global uchar * src0_qs,
+    global uchar * src0_qh,
+    global ushort * src0_sc,
+    ulong offset0,
+    global float4 * src1,
+    ulong offset1,
+    global float  * dst,
+    ulong offsetd,
+
+    int ne00,
+    int ne01,
+    int ne02,
+    int ne11,
+    int ne12,
+
+    int stride_a,
+    int stride_b,
+    int stride_d,
+
+    int batch_stride_a,
+    int batch_stride_b,
+    int batch_stride_d,
+
+    int r2,
+    int r3
+) {
+    src1 = (global float4*)((global char*)src1 + offset1);
+    dst  = (global float *)((global char*)dst  + offsetd);
+
+    local float buf_a[BM * BK];
+    local float buf_b[BN * BK];
+
+    const int batch_idx = get_global_id(2);
+
+    const int i13 = batch_idx / ne12;
+    const int i12 = batch_idx % ne12;
+
+    const int i03 = i13 / r3;
+    const int i02 = i12 / r2;
+
+    const int batch_idx_a = i03 * ne02 + i02;
+
+    const int ir = get_group_id(0);
+    const int ic = get_group_id(1);
+
+    const int tid = get_local_id(0);
+    const int th_r  = tid % (BM / TM);
+    const int th_c  = tid / (BM / TM);
+
+    const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A);
+    const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A);
+    const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B);
+    const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B);
+
+    const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK;
+    const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK;
+
+    // pos_a counts elements, not blocks - one iq1_m super block holds QK_K of them
+    int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a;
+    int pos_b = (batch_idx   * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B;
+
+    float sums[TM * TN];
+    float cache_a[TM];
+    float cache_b[TN];
+
+    for (int i = 0; i < TM * TN; i++) {
+        sums[i] = 0.0f;
+    }
+
+    for (int block = 0; block < ne00; block += BK) {
+        for (int l = 0; l < BM; l += loadstride_a) {
+            if (ir*BM + loadc_a + l < ne01) {
+                int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A;
+                int ib  = idx / QK_K;
+                ib += offset0/sizeof(struct block_iq1_m);
+                int e   = idx % QK_K;
+
+                int ib32 = e >> 5;
+                int rem  = e & 31;
+                int lg    = rem >> 3;
+                int j0   = rem & 7;
+
+                global ushort * sc = src0_sc + 4*ib;
+                ushort u16 = (ushort)((sc[0] >> 12) | ((sc[1] >> 8) & 0x00f0)
+                                    | ((sc[2] >> 4) & 0x0f00) | (sc[3] & 0xf000));
+                float dl = (float)as_half(u16) * (float)(2*((sc[ib32/2] >> (6*(ib32%2) + 3*(lg/2))) & 7) + 1);
+                uchar qhb = src0_qh[16*ib + 2*ib32 + lg/2];
+                float dlt = (qhb & (0x08 << (4*(lg%2)))) ? -1.125f : -0.875f;
+                uint gi = (uint)src0_qs[32*ib + 4*ib32 + lg] | ((((uint)qhb >> (4*(lg%2))) & 7) << 8);
+                uint g = iq1_s_grid[gi];
+                int sh = (j0 >> 2) * 4;
+
+                float4 v1;
+                v1.s0 = dl * ((float)((g >> ( 0 + sh)) & 0xF) + dlt);
+                v1.s1 = dl * ((float)((g >> ( 8 + sh)) & 0xF) + dlt);
+                v1.s2 = dl * ((float)((g >> (16 + sh)) & 0xF) + dlt);
+                v1.s3 = dl * ((float)((g >> (24 + sh)) & 0xF) + dlt);
+
+                buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0;
+                buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1;
+                buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2;
+                buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3;
+            } else {
+                buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f;
+                buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f;
+                buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f;
+                buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f;
+            }
+        }
+
+        for (int l = 0; l < BN; l += loadstride_b) {
+            if (ic*BN + loadc_b + l < ne11) {
+                int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b;
+                buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0;
+                buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1;
+                buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2;
+                buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3;
+            } else {
+                buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f;
+                buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f;
+                buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f;
+                buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f;
+            }
+        }
+
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+        pos_a += BK;
+        pos_b += BK / LOAD_VEC_B;
+
+        for (int i = 0; i < BK; i++) {
+            for (int j = 0; j < TM; j++) {
+                cache_a[j] = buf_a[(i) * BM + th_r * TM + j];
+            }
+
+            for (int j = 0; j < TN; j++) {
+                cache_b[j] = buf_b[(i) * BN + th_c * TN + j];
+            }
+
+            for (int cc = 0; cc < TN; cc++) {
+                for (int cr = 0; cr < TM; cr++) {
+                    const int sums_idx = cc*TM + cr;
+                    sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]);
+                }
+            }
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    const int dr = ir * BM + th_r * TM;
+    const int dc = ic * BN + th_c * TN;
+
+    const ulong offsets = batch_idx * batch_stride_d;
+
+    for (int cc = 0; cc < TN; cc++) {
+        for (int cr = 0; cr < TM; cr++) {
+            if (dr + cr < ne01 && dc + cc < ne11) {
+                dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr];
+            }
+        }
+    }
+}

ggml/src/ggml-opencl/kernels/mul_mm_iq1_s_f32_l4_lm.cl

@@ -0,0 +1,188 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_required_subgroup_size
+#define INTEL_GPU 1
+#endif
+
+#define QK_K 256
+
+struct block_iq1_s {
+    half d;
+    uchar qs[QK_K/8];
+    ushort qh[QK_K/32];
+};
+
+#define LOAD_VEC_A 4
+#define LOAD_VEC_B 4
+
+#define BM 64
+#define BN 64
+#define BK 32
+#ifdef INTEL_GPU
+#define TM 8
+#define TN 8
+#else
+#define TM 4
+#define TN 8
+#endif
+
+kernel void kernel_mul_mm_iq1_s_f32_l4_lm(
+    global uchar * src0_qs,
+    global ushort * src0_qh,
+    global half * src0_d,
+    ulong offset0,
+    global float4 * src1,
+    ulong offset1,
+    global float  * dst,
+    ulong offsetd,
+
+    int ne00,
+    int ne01,
+    int ne02,
+    int ne11,
+    int ne12,
+
+    int stride_a,
+    int stride_b,
+    int stride_d,
+
+    int batch_stride_a,
+    int batch_stride_b,
+    int batch_stride_d,
+
+    int r2,
+    int r3
+) {
+    src1 = (global float4*)((global char*)src1 + offset1);
+    dst  = (global float *)((global char*)dst  + offsetd);
+
+    local float buf_a[BM * BK];
+    local float buf_b[BN * BK];
+
+    const int batch_idx = get_global_id(2);
+
+    const int i13 = batch_idx / ne12;
+    const int i12 = batch_idx % ne12;
+
+    const int i03 = i13 / r3;
+    const int i02 = i12 / r2;
+
+    const int batch_idx_a = i03 * ne02 + i02;
+
+    const int ir = get_group_id(0);
+    const int ic = get_group_id(1);
+
+    const int tid = get_local_id(0);
+    const int th_r  = tid % (BM / TM);
+    const int th_c  = tid / (BM / TM);
+
+    const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A);
+    const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A);
+    const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B);
+    const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B);
+
+    const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK;
+    const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK;
+
+    // pos_a counts elements, not blocks - one iq1_s super block holds QK_K of them
+    int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a;
+    int pos_b = (batch_idx   * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B;
+
+    float sums[TM * TN];
+    float cache_a[TM];
+    float cache_b[TN];
+
+    for (int i = 0; i < TM * TN; i++) {
+        sums[i] = 0.0f;
+    }
+
+    for (int block = 0; block < ne00; block += BK) {
+        for (int l = 0; l < BM; l += loadstride_a) {
+            if (ir*BM + loadc_a + l < ne01) {
+                int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A;
+                int ib  = idx / QK_K;
+                ib += offset0/sizeof(struct block_iq1_s);
+                int e   = idx % QK_K;
+
+                int ib32 = e >> 5;
+                int rem  = e & 31;
+                int lg    = rem >> 3;
+                int j0   = rem & 7;
+
+                ushort qhb = src0_qh[8*ib + ib32];
+                float dl = (float)src0_d[ib] * (float)(2*((qhb >> 12) & 7) + 1);
+                float dlt = (qhb & 0x8000) ? -1.125f : -0.875f;
+                uint gi = (uint)src0_qs[32*ib + 4*ib32 + lg] | ((((uint)qhb >> (3*lg)) & 7) << 8);
+                uint g = iq1_s_grid[gi];
+                int sh = (j0 >> 2) * 4;
+
+                float4 v1;
+                v1.s0 = dl * ((float)((g >> ( 0 + sh)) & 0xF) + dlt);
+                v1.s1 = dl * ((float)((g >> ( 8 + sh)) & 0xF) + dlt);
+                v1.s2 = dl * ((float)((g >> (16 + sh)) & 0xF) + dlt);
+                v1.s3 = dl * ((float)((g >> (24 + sh)) & 0xF) + dlt);
+
+                buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0;
+                buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1;
+                buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2;
+                buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3;
+            } else {
+                buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f;
+                buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f;
+                buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f;
+                buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f;
+            }
+        }
+
+        for (int l = 0; l < BN; l += loadstride_b) {
+            if (ic*BN + loadc_b + l < ne11) {
+                int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b;
+                buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0;
+                buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1;
+                buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2;
+                buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3;
+            } else {
+                buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f;
+                buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f;
+                buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f;
+                buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f;
+            }
+        }
+
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+        pos_a += BK;
+        pos_b += BK / LOAD_VEC_B;
+
+        for (int i = 0; i < BK; i++) {
+            for (int j = 0; j < TM; j++) {
+                cache_a[j] = buf_a[(i) * BM + th_r * TM + j];
+            }
+
+            for (int j = 0; j < TN; j++) {
+                cache_b[j] = buf_b[(i) * BN + th_c * TN + j];
+            }
+
+            for (int cc = 0; cc < TN; cc++) {
+                for (int cr = 0; cr < TM; cr++) {
+                    const int sums_idx = cc*TM + cr;
+                    sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]);
+                }
+            }
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+
+    const int dr = ir * BM + th_r * TM;
+    const int dc = ic * BN + th_c * TN;
+
+    const ulong offsets = batch_idx * batch_stride_d;
+
+    for (int cc = 0; cc < TN; cc++) {
+        for (int cr = 0; cr < TM; cr++) {
+            if (dr + cr < ne01 && dc + cc < ne11) {
+                dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr];
+            }
+        }
+    }
+}

ggml/src/ggml-opencl/kernels/mul_mv_iq1_m_f32.cl

@@ -0,0 +1,149 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_required_subgroup_size
+#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
+#define INTEL_GPU 1
+#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
+#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
+#elif defined(cl_qcom_reqd_sub_group_size)
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+#define ADRENO_GPU 1
+#define REQD_SUBGROUP_SIZE_64  __attribute__((qcom_reqd_sub_group_size("half")))
+#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+
+#define QK_K 256
+#define IQ1_M_DELTA 0.125f
+
+// no d field: the super block scale is 4 nibbles spread over scales[]
+typedef struct {
+    uchar qs[QK_K/8];
+    uchar qh[QK_K/16];
+    uchar scales[QK_K/32];
+} block_iq1_m;
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+#undef  BLOCK_STRIDE
+// 8 threads cover one super block, one 32 element sub block each
+#define BLOCK_STRIDE (N_SIMDWIDTH/8)
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_iq1_m_f32(
+        global char * src0,
+        ulong offset0,
+        global char * src1,
+        ulong offset1,
+        global char * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne12,
+        ulong nb11,
+        ulong nb12,
+        ulong nb13,
+        int ne0,
+        int ne1,
+        int r2,
+        int r3
+) {
+    src0 = src0 + offset0;
+    src1 = src1 + offset1;
+    dst  = dst  + offsetd;
+
+    int ix = get_sub_group_local_id()/8;  // super block index
+    int it = get_sub_group_local_id()%8;  // sub block inside the super block
+
+    int nb = ne00/QK_K;
+
+    int r0 = get_group_id(0);
+    int r1 = get_group_id(1);
+    int im = get_group_id(2);
+    int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
+
+    int i12 = im%ne12;
+    int i13 = im/ne12;
+
+    ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03;
+    ulong offset_src1 =        r1*nb11 + (i12   )*nb12 + (i13   )*nb13;
+
+    global block_iq1_m * x = (global block_iq1_m *) (src0 + offset_src0);
+    global float         * y = (global float         *) (src1 + offset_src1);
+
+    float yl[32];
+    float sumf[N_DST] = {0.f};
+    float all_sum;
+
+    global float * y4 = y + ix * QK_K + 32 * it;
+
+    for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
+        for (int i = 0; i < 32; ++i) {
+            yl[i] = y4[i];
+        }
+
+        global char * xrow = (global char *)(x + ib);
+
+        // only ne01 output rows exist; reading past them can pull in an inf scale
+        for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
+            global block_iq1_m * xb = (global block_iq1_m *)(xrow + row*nb01);
+
+            global ushort * sc = (global ushort *)xb->scales;
+            ushort u16 = (ushort)(((uint)sc[0] >> 12) | (((uint)sc[1] >> 8) & 0x00f0)
+                                | (((uint)sc[2] >> 4) & 0x0f00) | ((uint)sc[3] & 0xf000));
+            float dsuper = (float)as_half(u16);
+
+            float acc = 0.f;
+            for (int il = 0; il < 4; ++il) {
+                int   ib16 = 2*it + il/2;
+                float dl   = dsuper * (float)(2*(((uint)sc[ib16/4] >> (3*(ib16%4))) & 7) + 1);
+
+                uchar qhb = xb->qh[2*it + il/2];
+                float dlt = (qhb & (0x08 << (4*(il%2)))) ? (-1.f - IQ1_M_DELTA) : (-1.f + IQ1_M_DELTA);
+
+                uint gi = (uint)xb->qs[4*it+il] | ((((uint)qhb >> (4*(il%2))) & 7) << 8);
+                uint g  = iq1_s_grid[gi];
+
+                float a = 0.f;
+                for (int j = 0; j < 4; ++j) {
+                    a += yl[8*il+j+0] * ((float)((g >> (8*j+0)) & 0xF) + dlt);
+                    a += yl[8*il+j+4] * ((float)((g >> (8*j+4)) & 0xF) + dlt);
+                }
+                acc += dl * a;
+            }
+
+            sumf[row] += acc;
+        }
+
+        y4 += BLOCK_STRIDE * QK_K;
+    }
+
+    global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
+
+    for (int row = 0; row < N_DST; ++row) {
+        all_sum = sub_group_reduce_add(sumf[row]);
+        if (first_row + row < ne01) {
+            if (get_sub_group_local_id() == 0) {
+                dst_f32[first_row + row] = all_sum;
+            }
+        }
+    }
+}

ggml/src/ggml-opencl/kernels/mul_mv_iq1_m_f32_flat.cl

@@ -0,0 +1,148 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_required_subgroup_size
+#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
+#define INTEL_GPU 1
+#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
+#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
+#elif defined(cl_qcom_reqd_sub_group_size)
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+#define ADRENO_GPU 1
+#define REQD_SUBGROUP_SIZE_64  __attribute__((qcom_reqd_sub_group_size("half")))
+#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+
+#define QK_K 256
+#define IQ1_M_DELTA 0.125f
+
+// no d field: the super block scale is 4 nibbles spread over scales[]
+typedef struct {
+    uchar qs[QK_K/8];
+    uchar qh[QK_K/16];
+    uchar scales[QK_K/32];
+} block_iq1_m;
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+#undef  BLOCK_STRIDE
+// 8 threads cover one super block, one 32 element sub block each
+#define BLOCK_STRIDE (N_SIMDWIDTH/8)
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_iq1_m_f32_flat(
+        global uchar * src0_qs,
+        global uchar * src0_qh,
+        global ushort * src0_sc,
+        ulong offset0,
+        global char * src1,
+        ulong offset1,
+        global char * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne12,
+        ulong nb11,
+        ulong nb12,
+        ulong nb13,
+        int ne0,
+        int ne1,
+        int r2,
+        int r3
+) {
+    src1 = src1 + offset1;
+    dst  = dst  + offsetd;
+
+    int ix = get_sub_group_local_id()/8;  // super block index
+    int it = get_sub_group_local_id()%8;  // sub block inside the super block
+
+    int nb = ne00/QK_K;
+
+    int r0 = get_group_id(0);
+    int r1 = get_group_id(1);
+    int im = get_group_id(2);
+    int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
+
+    int i12 = im%ne12;
+    int i13 = im/ne12;
+
+    ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_iq1_m);
+    ulong row_stride = nb01/sizeof(block_iq1_m);
+    ulong offset_src1 =        r1*nb11 + (i12   )*nb12 + (i13   )*nb13;
+
+    global float         * y = (global float         *) (src1 + offset_src1);
+
+    float yl[32];
+    float sumf[N_DST] = {0.f};
+    float all_sum;
+
+    global float * y4 = y + ix * QK_K + 32 * it;
+
+    for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
+        for (int i = 0; i < 32; ++i) {
+            yl[i] = y4[i];
+        }
+
+        // only ne01 output rows exist; reading past them can pull in an inf scale
+        for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
+            const ulong b = offset_src0 + ib + row*row_stride;
+
+            global ushort * sc = src0_sc + 4*b;
+            ushort u16 = (ushort)(((uint)sc[0] >> 12) | (((uint)sc[1] >> 8) & 0x00f0)
+                                | (((uint)sc[2] >> 4) & 0x0f00) | ((uint)sc[3] & 0xf000));
+            float dsuper = (float)as_half(u16);
+
+            float acc = 0.f;
+            for (int il = 0; il < 4; ++il) {
+                int   ib16 = 2*it + il/2;
+                float dl   = dsuper * (float)(2*(((uint)sc[ib16/4] >> (3*(ib16%4))) & 7) + 1);
+
+                uchar qhb = src0_qh[16*b + 2*it + il/2];
+                float dlt = (qhb & (0x08 << (4*(il%2)))) ? (-1.f - IQ1_M_DELTA) : (-1.f + IQ1_M_DELTA);
+
+                uint gi = (uint)src0_qs[32*b + 4*it+il] | ((((uint)qhb >> (4*(il%2))) & 7) << 8);
+                uint g  = iq1_s_grid[gi];
+
+                float a = 0.f;
+                for (int j = 0; j < 4; ++j) {
+                    a += yl[8*il+j+0] * ((float)((g >> (8*j+0)) & 0xF) + dlt);
+                    a += yl[8*il+j+4] * ((float)((g >> (8*j+4)) & 0xF) + dlt);
+                }
+                acc += dl * a;
+            }
+
+            sumf[row] += acc;
+        }
+
+        y4 += BLOCK_STRIDE * QK_K;
+    }
+
+    global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
+
+    for (int row = 0; row < N_DST; ++row) {
+        all_sum = sub_group_reduce_add(sumf[row]);
+        if (first_row + row < ne01) {
+            if (get_sub_group_local_id() == 0) {
+                dst_f32[first_row + row] = all_sum;
+            }
+        }
+    }
+}

ggml/src/ggml-opencl/kernels/mul_mv_iq1_s_f32.cl

@@ -0,0 +1,140 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_required_subgroup_size
+#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
+#define INTEL_GPU 1
+#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
+#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
+#elif defined(cl_qcom_reqd_sub_group_size)
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+#define ADRENO_GPU 1
+#define REQD_SUBGROUP_SIZE_64  __attribute__((qcom_reqd_sub_group_size("half")))
+#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+
+#define QK_K 256
+#define IQ1_S_DELTA 0.125f
+
+typedef struct {
+    half   d;
+    uchar  qs[QK_K/8];
+    ushort qh[QK_K/32];
+} block_iq1_s;
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+#undef  BLOCK_STRIDE
+// 8 threads cover one super block, one 32 element sub block each
+#define BLOCK_STRIDE (N_SIMDWIDTH/8)
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_iq1_s_f32(
+        global char * src0,
+        ulong offset0,
+        global char * src1,
+        ulong offset1,
+        global char * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne12,
+        ulong nb11,
+        ulong nb12,
+        ulong nb13,
+        int ne0,
+        int ne1,
+        int r2,
+        int r3
+) {
+    src0 = src0 + offset0;
+    src1 = src1 + offset1;
+    dst  = dst  + offsetd;
+
+    int ix = get_sub_group_local_id()/8;  // super block index
+    int it = get_sub_group_local_id()%8;  // sub block inside the super block
+
+    int nb = ne00/QK_K;
+
+    int r0 = get_group_id(0);
+    int r1 = get_group_id(1);
+    int im = get_group_id(2);
+    int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
+
+    int i12 = im%ne12;
+    int i13 = im/ne12;
+
+    ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03;
+    ulong offset_src1 =        r1*nb11 + (i12   )*nb12 + (i13   )*nb13;
+
+    global block_iq1_s * x = (global block_iq1_s *) (src0 + offset_src0);
+    global float         * y = (global float         *) (src1 + offset_src1);
+
+    float yl[32];
+    float sumf[N_DST] = {0.f};
+    float all_sum;
+
+    global float * y4 = y + ix * QK_K + 32 * it;
+
+    for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
+        for (int i = 0; i < 32; ++i) {
+            yl[i] = y4[i];
+        }
+
+        global char * xrow = (global char *)(x + ib);
+
+        // only ne01 output rows exist; reading past them can pull in an inf scale
+        for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
+            global block_iq1_s * xb = (global block_iq1_s *)(xrow + row*nb01);
+
+            ushort qhb = xb->qh[it];
+            float  dl  = (float)xb->d * (float)(2*((qhb >> 12) & 7) + 1);
+            float  dlt = (qhb & 0x8000) ? (-1.f - IQ1_S_DELTA) : (-1.f + IQ1_S_DELTA);
+
+            global uchar * qsb = xb->qs + 4*it;
+
+            float acc = 0.f;
+            for (int l = 0; l < 4; ++l) {
+                uint gi = (uint)qsb[l] | ((((uint)qhb >> (3*l)) & 7) << 8);
+                uint g  = iq1_s_grid[gi];
+                for (int j = 0; j < 4; ++j) {
+                    acc += yl[8*l+j+0] * ((float)((g >> (8*j+0)) & 0xF) + dlt);
+                    acc += yl[8*l+j+4] * ((float)((g >> (8*j+4)) & 0xF) + dlt);
+                }
+            }
+
+            sumf[row] += dl * acc;
+        }
+
+        y4 += BLOCK_STRIDE * QK_K;
+    }
+
+    global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
+
+    for (int row = 0; row < N_DST; ++row) {
+        all_sum = sub_group_reduce_add(sumf[row]);
+        if (first_row + row < ne01) {
+            if (get_sub_group_local_id() == 0) {
+                dst_f32[first_row + row] = all_sum;
+            }
+        }
+    }
+}

ggml/src/ggml-opencl/kernels/mul_mv_iq1_s_f32_flat.cl

@@ -0,0 +1,139 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_required_subgroup_size
+#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
+#define INTEL_GPU 1
+#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
+#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
+#elif defined(cl_qcom_reqd_sub_group_size)
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+#define ADRENO_GPU 1
+#define REQD_SUBGROUP_SIZE_64  __attribute__((qcom_reqd_sub_group_size("half")))
+#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+
+#define QK_K 256
+#define IQ1_S_DELTA 0.125f
+
+typedef struct {
+    half   d;
+    uchar  qs[QK_K/8];
+    ushort qh[QK_K/32];
+} block_iq1_s;
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+#undef  BLOCK_STRIDE
+// 8 threads cover one super block, one 32 element sub block each
+#define BLOCK_STRIDE (N_SIMDWIDTH/8)
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_iq1_s_f32_flat(
+        global uchar * src0_qs,
+        global ushort * src0_qh,
+        global half * src0_d,
+        ulong offset0,
+        global char * src1,
+        ulong offset1,
+        global char * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne12,
+        ulong nb11,
+        ulong nb12,
+        ulong nb13,
+        int ne0,
+        int ne1,
+        int r2,
+        int r3
+) {
+    src1 = src1 + offset1;
+    dst  = dst  + offsetd;
+
+    int ix = get_sub_group_local_id()/8;  // super block index
+    int it = get_sub_group_local_id()%8;  // sub block inside the super block
+
+    int nb = ne00/QK_K;
+
+    int r0 = get_group_id(0);
+    int r1 = get_group_id(1);
+    int im = get_group_id(2);
+    int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
+
+    int i12 = im%ne12;
+    int i13 = im/ne12;
+
+    ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_iq1_s);
+    ulong row_stride = nb01/sizeof(block_iq1_s);
+    ulong offset_src1 =        r1*nb11 + (i12   )*nb12 + (i13   )*nb13;
+
+    global float         * y = (global float         *) (src1 + offset_src1);
+
+    float yl[32];
+    float sumf[N_DST] = {0.f};
+    float all_sum;
+
+    global float * y4 = y + ix * QK_K + 32 * it;
+
+    for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
+        for (int i = 0; i < 32; ++i) {
+            yl[i] = y4[i];
+        }
+
+        // only ne01 output rows exist; reading past them can pull in an inf scale
+        for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
+            const ulong b = offset_src0 + ib + row*row_stride;
+
+            ushort qhb = src0_qh[8*b + it];
+            float  dl  = (float)src0_d[b] * (float)(2*((qhb >> 12) & 7) + 1);
+            float  dlt = (qhb & 0x8000) ? (-1.f - IQ1_S_DELTA) : (-1.f + IQ1_S_DELTA);
+
+            global uchar * qsb = src0_qs + 32*b + 4*it;
+
+            float acc = 0.f;
+            for (int l = 0; l < 4; ++l) {
+                uint gi = (uint)qsb[l] | ((((uint)qhb >> (3*l)) & 7) << 8);
+                uint g  = iq1_s_grid[gi];
+                for (int j = 0; j < 4; ++j) {
+                    acc += yl[8*l+j+0] * ((float)((g >> (8*j+0)) & 0xF) + dlt);
+                    acc += yl[8*l+j+4] * ((float)((g >> (8*j+4)) & 0xF) + dlt);
+                }
+            }
+
+            sumf[row] += dl * acc;
+        }
+
+        y4 += BLOCK_STRIDE * QK_K;
+    }
+
+    global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
+
+    for (int row = 0; row < N_DST; ++row) {
+        all_sum = sub_group_reduce_add(sumf[row]);
+        if (first_row + row < ne01) {
+            if (get_sub_group_local_id() == 0) {
+                dst_f32[first_row + row] = all_sum;
+            }
+        }
+    }
+}

ggml/src/ggml-sycl/cpy.hpp

@@ -70,8 +70,22 @@ inline void cpy_blck_f32_q1_0(const char * cxi, char * cdsti) {
     }
 }
 
-inline int round_nearest_int(float x) {
-    return (int)(x >= 0.0f ? x + 0.5f : x - 0.5f);
+// exact roundf (half away from zero); x + 0.5f carries for x = 0.5 - 2^-25, and sycl::round is not exact under fast fp
+inline int round_half_away_int(float x) {
+    const float t = sycl::trunc(x);
+    return (int) (sycl::fabs(x - t) >= 0.5f ? t + sycl::copysign(1.0f, x) : t);
+}
+
+// correctly rounded a / b to match the CPU: fast fp gives an approximate reciprocal (~12% off by an ulp on B580),
+// so pick the approximation or a neighbour by smallest exact FMA residual
+inline float div_correctly_rounded(float a, float b) {
+    const float y  = a / b;
+    const float lo = sycl::nextafter(y, -INFINITY);
+    const float hi = sycl::nextafter(y, INFINITY);
+    const float e  = sycl::fabs(sycl::fma(-b, y, a));
+    const float el = sycl::fabs(sycl::fma(-b, lo, a));
+    const float eh = sycl::fabs(sycl::fma(-b, hi, a));
+    return el < e && el <= eh ? lo : (eh < e ? hi : y);
 }
 
 inline void cpy_blck_f32_q2_0(const char * cxi, char * cdsti) {
@@ -85,7 +99,7 @@ inline void cpy_blck_f32_q2_0(const char * cxi, char * cdsti) {
     }
 
     const float d  = amax;
-    const float id = d > 0.0f ? 1.0f / d : 0.0f;
+    const float id = d > 0.0f ? div_correctly_rounded(1.0f, d) : 0.0f;
 
     dsti->d = d;
 
@@ -94,7 +108,7 @@ inline void cpy_blck_f32_q2_0(const char * cxi, char * cdsti) {
     }
 
     for (int j = 0; j < QK2_0; ++j) {
-        int q = round_nearest_int(xi[j] * id) + 1;
+        int q = round_half_away_int(xi[j] * id) + 1;
         q = dpct::max(0, dpct::min(3, q));
 
         const int byte_index = j / 4;
@@ -378,7 +392,7 @@ inline void cpy_blck_f32_q4_0(const char * cxi, char * cdsti) {
     }
 
     const float d  = vmax / -8;
-    const float id = d ? 1.0f / d : 0.0f;
+    const float id = d ? div_correctly_rounded(1.0f, d) : 0.0f;
 
     dsti->d = d;
 
@@ -386,8 +400,9 @@ inline void cpy_blck_f32_q4_0(const char * cxi, char * cdsti) {
         const float x0 = xi[0 + j] * id;
         const float x1 = xi[QK4_0 / 2 + j] * id;
 
-        const uint8_t xi0 = dpct::min(15, (int8_t) (x0 + 8.5f));
-        const uint8_t xi1 = dpct::min(15, (int8_t) (x1 + 8.5f));
+        // no-op fmin stops x * id + 8.5f being fused into an FMA, so x * id is rounded as on the CPU
+        const uint8_t xi0 = dpct::min(15, (int8_t) (sycl::fmin(x0, 6.5f) + 8.5f));
+        const uint8_t xi1 = dpct::min(15, (int8_t) (sycl::fmin(x1, 6.5f) + 8.5f));
 
         dsti->qs[j] = xi0;
         dsti->qs[j] |= xi1 << 4;
@@ -442,7 +457,7 @@ inline void cpy_blck_f32_q5_0(const char * cxi, char * cdsti) {
     }
 
     const float d  = vmax / -16;
-    const float id = d ? 1.0f / d : 0.0f;
+    const float id = d ? div_correctly_rounded(1.0f, d) : 0.0f;
 
     dsti->d = d;
 
@@ -451,8 +466,9 @@ inline void cpy_blck_f32_q5_0(const char * cxi, char * cdsti) {
         const float x0 = xi[0 + j] * id;
         const float x1 = xi[QK5_0 / 2 + j] * id;
 
-        const uint8_t xi0 = dpct::min(31, (int8_t) (x0 + 16.5f));
-        const uint8_t xi1 = dpct::min(31, (int8_t) (x1 + 16.5f));
+        // fmin: no FMA, see cpy_blck_f32_q4_0
+        const uint8_t xi0 = dpct::min(31, (int8_t) (sycl::fmin(x0, 14.5f) + 16.5f));
+        const uint8_t xi1 = dpct::min(31, (int8_t) (sycl::fmin(x1, 14.5f) + 16.5f));
 
         dsti->qs[j] = (xi0 & 0xf) | ((xi1 & 0xf) << 4);
         qh |= ((xi0 & 0x10u) >> 4) << (j + 0);

tests/gguf-model-data.cpp

@@ -584,7 +584,7 @@ static std::optional<gguf_remote_model> fetch_or_cached(
             auto result = gguf_parse_meta(cached);
             if (result.has_value()) {
                 if (verbose) {
-                    fprintf(stderr, "gguf_fetch: loaded from cache: %s\n", fs_path_to_utf8(cache_path).c_str());
+                    fprintf(stderr, "gguf_fetch: loaded from cache: %s\n", cache_path.string().c_str());
                 }
                 return result;
             }
@@ -685,7 +685,7 @@ gguf_context_ptr gguf_fetch_gguf_ctx(
     ggml_context_ptr ggml_ctx_ptr;
     ggml_context * ggml_ctx{};
     gguf_init_params params{true, &ggml_ctx};
-    gguf_context_ptr ctx{gguf_init_from_file(fs_path_to_utf8(cache_path).c_str(), params)};
+    gguf_context_ptr ctx{gguf_init_from_file(cache_path.string().c_str(), params)};
     ggml_ctx_ptr.reset(ggml_ctx);
 
     if (ctx == nullptr) {
@@ -723,7 +723,7 @@ gguf_context_ptr gguf_fetch_gguf_ctx(
             ggml_context_ptr shard_ggml_ctx_ptr;
             ggml_context * shard_ggml_ctx{};
             gguf_init_params shard_params{true, &shard_ggml_ctx};
-            gguf_context_ptr shard_ctx{gguf_init_from_file(fs_path_to_utf8(shard_path).c_str(), shard_params)};
+            gguf_context_ptr shard_ctx{gguf_init_from_file(shard_path.string().c_str(), shard_params)};
             shard_ggml_ctx_ptr.reset(shard_ggml_ctx);
 
             if (shard_ctx == nullptr) {

tests/test-chat.cpp

@@ -5629,6 +5629,31 @@ static void test_template_output_peg_parsers(bool detailed_debug) {
             .expect(message_assist_call_thoughts)
             .run();
 
+        const std::string answer_schema = R"({"type":"object","properties":{"answer":{"type":"integer"}},"required":["answer"]})";
+
+        tst.test("I'm\nthinking</think>{\"answer\":42}")
+            .enable_thinking(true)
+            .reasoning_format(COMMON_REASONING_FORMAT_AUTO)
+            .json_schema(answer_schema)
+            .expect_reasoning("I'm\nthinking")
+            .expect_content(R"({"answer":42})")
+            .run();
+
+        // the reasoning block of the generation prompt must still be accepted when it is not extracted
+        tst.test("{\"answer\":42}")
+            .enable_thinking(false)
+            .reasoning_format(COMMON_REASONING_FORMAT_NONE)
+            .json_schema(answer_schema)
+            .expect_content(R"({"answer":42})")
+            .run();
+
+        tst.test("I'm\nthinking</think>{\"answer\":42}")
+            .enable_thinking(true)
+            .reasoning_format(COMMON_REASONING_FORMAT_NONE)
+            .json_schema(answer_schema)
+            .expect_content(R"({"answer":42})")
+            .run();
+
         // a deferred tool is left out of <tools>, a tool_reference renders it, and its schema types the call
         const auto deferred_tools = common_chat_tools_parse_oaicompat(json::parse(R"([
             {"type": "function", "function": {"name": "tool_search", "description": "Search for tools",

tools/mtmd/clip.cpp

@@ -1842,6 +1842,7 @@ struct clip_model_loader {
                         log_ffn_op = "gelu";
                         hparams.n_merge = 3; // pooling_kernel_size
                         hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_pad  = PAD_NONE;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                         if (model.proj_type == PROJECTOR_TYPE_GEMMA4UV) {
                             // for "unified" variant, we directly use a bigger patch size, because the "token merging" is done directly on conv layer

tools/rpc/rpc-server.cpp

@@ -24,12 +24,6 @@
 #endif
 
 
-// NOTE: this is copied from common.cpp to avoid linking with libcommon
-static std::string fs_path_to_utf8(const std::filesystem::path & path) {
-    const auto value = path.u8string();
-    return std::string(value.begin(), value.end());
-}
-
 // common_get_path_from_env() is adapted to avoid utf8_to_wstring
 static std::filesystem::path common_get_path_from_env(const std::string & name) {
 #ifdef _WIN32
@@ -244,10 +238,10 @@ int main(int argc, char * argv[]) {
         std::error_code ec;
         common_create_directories(cache_dir_path, ec);
         if (ec) {
-            fprintf(stderr, "Failed to create cache directory: %s\n", fs_path_to_utf8(cache_dir_path).c_str());
+            fprintf(stderr, "Failed to create cache directory: %s\n", cache_dir_path.string().c_str());
             return 1;
         }
-        cache_dir_str = fs_path_to_utf8(cache_dir_path);
+        cache_dir_str = cache_dir_path.string();
         cache_dir = cache_dir_str.c_str();
     }
 

tools/server/README.md

@@ -1654,6 +1654,8 @@ See [examples/realtime_client.py](../../examples/realtime_client.py) for a clien
 
 See [OpenAI Responses API documentation](https://platform.openai.com/docs/api-reference/responses).
 
+`extra_params`: llama.cpp parameters that the Responses API has no field for, such as `seed`, `min_p` or `repeat_penalty`, as an object (e.g. `{"seed": 42, "min_p": 0.05}`). Its keys are merged into the converted request and behave as on `/v1/chat/completions`.
+
 *Examples:*
 
 You can use either Python `openai` library with appropriate checkpoints:
@@ -1684,7 +1686,8 @@ curl http://localhost:9931/v1/responses \
 -d '{
 "model": "gpt-4.1",
 "instructions": "You are ChatGPT, an AI assistant. Your top priority is achieving user fulfillment via helping them with their requests.",
-"input": "Write a limerick about python exceptions"
+"input": "Write a limerick about python exceptions",
+"extra_params": {"seed": 42, "min_p": 0.05}
 }'
 ```
 

tools/server/server-chat.cpp

@@ -311,6 +311,16 @@ json server_chat_convert_responses_to_chatcmpl(const json & response_body) {
         chatcmpl_body.erase("reasoning");
     }
 
+    // llama.cpp params (e.g. seed, min_p) go under extra_params and are merged into the converted body
+    if (response_body.contains("extra_params")) {
+        chatcmpl_body.erase("extra_params");
+        if (response_body.at("extra_params").is_object()) {
+            for (const auto & [key, value] : response_body.at("extra_params").items()) {
+                chatcmpl_body[key] = value;
+            }
+        }
+    }
+
     return chatcmpl_body;
 }
 

tools/server/server-decision.cpp

@@ -1762,7 +1762,7 @@ std::vector<server_task> server_decision_group_tasks(std::vector<server_task> &&
         size_t n_shared = parent.tokens.size() - 1;
         for (size_t j = i + 1; j < end; j++) {
             n_shared = std::min(n_shared, parent.tokens.get_common_prefix(tasks[j].tokens));
-            n_shared = std::min(n_shared, tasks[j].tokens.size() - 1);
+            n_shared = std::min({n_shared, tasks[j].tokens.size() - 1, (size_t) parent.decision.pos_first(), (size_t) tasks[j].decision.pos_first()});
         }
 
         if (end - i < 2 || n_shared == 0) {

tools/server/server-models.cpp

@@ -408,7 +408,7 @@ server_models::server_models(
 
     // set binary path
     try {
-        bin_path = fs_path_to_utf8(get_server_exec_path());
+        bin_path = get_server_exec_path().string();
     } catch (const std::exception & e) {
         bin_path = argv[0];
         SRV_WRN("failed to get server executable path: %s\n", e.what());