metal : cap threadgroups size of set_rows

.devops/nix/package.nix

@@ -34,7 +34,6 @@
   rocmGpuTargets ? builtins.concatStringsSep ";" rocmPackages.clr.gpuTargets,
   enableCurl ? true,
   useVulkan ? false,
-  useRpc ? false,
   llamaVersion ? "0.0.0", # Arbitrary version, substituted by the flake
 
   # It's necessary to consistently use backendStdenv when building with CUDA support,
@@ -176,7 +175,6 @@ effectiveStdenv.mkDerivation (finalAttrs: {
       (cmakeBool "GGML_METAL" useMetalKit)
       (cmakeBool "GGML_VULKAN" useVulkan)
       (cmakeBool "GGML_STATIC" enableStatic)
-      (cmakeBool "GGML_RPC" useRpc)
     ]
     ++ optionals useCuda [
       (

convert_hf_to_gguf.py

@@ -7354,7 +7354,6 @@ class PLMModel(TextModel):
 @ModelBase.register("T5ForConditionalGeneration")
 @ModelBase.register("MT5ForConditionalGeneration")
 @ModelBase.register("UMT5ForConditionalGeneration")
-@ModelBase.register("UMT5Model")
 class T5Model(TextModel):
     model_arch = gguf.MODEL_ARCH.T5
 

ggml/src/ggml-cpu/CMakeLists.txt

@@ -126,36 +126,25 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
                 )
                 if (NOT ARM_MCPU_RESULT)
                     string(REGEX MATCH "-mcpu=[^ ']+" ARM_MCPU_FLAG "${ARM_MCPU}")
-                    string(REGEX MATCH "-march=[^ ']+" ARM_MARCH_FLAG "${ARM_MCPU}")
-
-                    # on some old GCC we need to read -march=
-                    if (ARM_MARCH_FLAG AND NOT "${ARM_MARCH_FLAG}" STREQUAL "-march=native")
-                        set(ARM_NATIVE_FLAG "${ARM_MARCH_FLAG}")
-                    elseif(ARM_MCPU_FLAG AND NOT "${ARM_MCPU_FLAG}" STREQUAL "-mcpu=native")
-                        set(ARM_NATIVE_FLAG "${ARM_MCPU_FLAG}")
-                    endif()
                 endif()
-
-                if ("${ARM_NATIVE_FLAG}" STREQUAL "")
-                    set(ARM_NATIVE_FLAG -mcpu=native)
-                    message(WARNING "ARM -march/-mcpu not found, -mcpu=native will be used")
-                else()
-                    message(STATUS "ARM detected flags: ${ARM_NATIVE_FLAG}")
+                if ("${ARM_MCPU_FLAG}" STREQUAL "")
+                    set(ARM_MCPU_FLAG -mcpu=native)
+                    message(STATUS "ARM -mcpu not found, -mcpu=native will be used")
                 endif()
 
                 include(CheckCXXSourceRuns)
 
                 function(check_arm_feature tag code)
                     set(CMAKE_REQUIRED_FLAGS_SAVE ${CMAKE_REQUIRED_FLAGS})
-                    set(CMAKE_REQUIRED_FLAGS "${ARM_NATIVE_FLAG}+${tag}")
+                    set(CMAKE_REQUIRED_FLAGS "${ARM_MCPU_FLAG}+${tag}")
                     check_cxx_source_runs("${code}" GGML_MACHINE_SUPPORTS_${tag})
                     if (GGML_MACHINE_SUPPORTS_${tag})
-                        set(ARM_NATIVE_FLAG_FIX "${ARM_NATIVE_FLAG_FIX}+${tag}" PARENT_SCOPE)
+                        set(ARM_MCPU_FLAG_FIX "${ARM_MCPU_FLAG_FIX}+${tag}" PARENT_SCOPE)
                     else()
-                        set(CMAKE_REQUIRED_FLAGS "${ARM_NATIVE_FLAG}+no${tag}")
+                        set(CMAKE_REQUIRED_FLAGS "${ARM_MCPU_FLAG}+no${tag}")
                         check_cxx_source_compiles("int main() { return 0; }" GGML_MACHINE_SUPPORTS_no${tag})
                         if (GGML_MACHINE_SUPPORTS_no${tag})
-                            set(ARM_NATIVE_FLAG_FIX "${ARM_NATIVE_FLAG_FIX}+no${tag}" PARENT_SCOPE)
+                            set(ARM_MCPU_FLAG_FIX "${ARM_MCPU_FLAG_FIX}+no${tag}" PARENT_SCOPE)
                         endif()
                     endif()
                     set(CMAKE_REQUIRED_FLAGS ${CMAKE_REQUIRED_FLAGS_SAVE})
@@ -166,7 +155,7 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
                 check_arm_feature(sve     "#include <arm_sve.h>\nint main()  { svfloat32_t _a, _b; volatile svfloat32_t _c = svadd_f32_z(svptrue_b8(), _a, _b); return 0; }")
                 check_arm_feature(sme     "#include <arm_sme.h>\n__arm_locally_streaming int main() { __asm__ volatile(\"smstart; smstop;\"); return 0; }")
 
-                list(APPEND ARCH_FLAGS "${ARM_NATIVE_FLAG}${ARM_NATIVE_FLAG_FIX}")
+                list(APPEND ARCH_FLAGS "${ARM_MCPU_FLAG}${ARM_MCPU_FLAG_FIX}")
             else()
                 if (GGML_CPU_ARM_ARCH)
                     list(APPEND ARCH_FLAGS -march=${GGML_CPU_ARM_ARCH})

ggml/src/ggml-cpu/ggml-cpu.c

@@ -1807,6 +1807,22 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
             {
                 ggml_compute_forward_cont(params, tensor);
             } break;
+        case GGML_OP_RESHAPE:
+            {
+                ggml_compute_forward_reshape(params, tensor);
+            } break;
+        case GGML_OP_VIEW:
+            {
+                ggml_compute_forward_view(params, tensor);
+            } break;
+        case GGML_OP_PERMUTE:
+            {
+                ggml_compute_forward_permute(params, tensor);
+            } break;
+        case GGML_OP_TRANSPOSE:
+            {
+                ggml_compute_forward_transpose(params, tensor);
+            } break;
         case GGML_OP_GET_ROWS:
             {
                 ggml_compute_forward_get_rows(params, tensor);
@@ -2026,22 +2042,6 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
             {
                 // nop
             } break;
-        case GGML_OP_RESHAPE:
-            {
-                // nop
-            } break;
-        case GGML_OP_PERMUTE:
-            {
-                // nop
-            } break;
-        case GGML_OP_VIEW:
-            {
-                // nop
-            } break;
-        case GGML_OP_TRANSPOSE:
-            {
-                // nop
-            } break;
         case GGML_OP_COUNT:
             {
                 GGML_ABORT("fatal error");
@@ -2884,11 +2884,6 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
     for (int node_n = 0; node_n < cgraph->n_nodes && atomic_load_explicit(&tp->abort, memory_order_relaxed) != node_n; node_n++) {
         struct ggml_tensor * node = cgraph->nodes[node_n];
 
-        if (ggml_op_is_empty(node->op)) {
-            // skip NOPs
-            continue;
-        }
-
         ggml_compute_forward(&params, node);
 
         if (state->ith == 0 && cplan->abort_callback &&

ggml/src/ggml-cpu/ops.cpp

@@ -4455,6 +4455,46 @@ void ggml_compute_forward_cont(
     ggml_compute_forward_dup(params, dst);
 }
 
+// ggml_compute_forward_reshape
+
+void ggml_compute_forward_reshape(
+        const ggml_compute_params * params,
+        ggml_tensor * dst) {
+    // NOP
+    GGML_UNUSED(params);
+    GGML_UNUSED(dst);
+}
+
+// ggml_compute_forward_view
+
+void ggml_compute_forward_view(
+        const ggml_compute_params * params,
+        ggml_tensor * dst) {
+    // NOP
+    GGML_UNUSED(params);
+    GGML_UNUSED(dst);
+}
+
+// ggml_compute_forward_permute
+
+void ggml_compute_forward_permute(
+        const ggml_compute_params * params,
+        ggml_tensor * dst) {
+    // NOP
+    GGML_UNUSED(params);
+    GGML_UNUSED(dst);
+}
+
+// ggml_compute_forward_transpose
+
+void ggml_compute_forward_transpose(
+        const ggml_compute_params * params,
+        ggml_tensor * dst) {
+    // NOP
+    GGML_UNUSED(params);
+    GGML_UNUSED(dst);
+}
+
 // ggml_compute_forward_get_rows
 
 static void ggml_compute_forward_get_rows_q(

ggml/src/ggml-cpu/ops.h

@@ -51,6 +51,10 @@ void ggml_compute_forward_scale(const struct ggml_compute_params * params, struc
 void ggml_compute_forward_set(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_cpy(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_cont(const struct ggml_compute_params * params, struct ggml_tensor * dst);
+void ggml_compute_forward_reshape(const struct ggml_compute_params * params, struct ggml_tensor * dst);
+void ggml_compute_forward_view(const struct ggml_compute_params * params, struct ggml_tensor * dst);
+void ggml_compute_forward_permute(const struct ggml_compute_params * params, struct ggml_tensor * dst);
+void ggml_compute_forward_transpose(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_get_rows(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_get_rows_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_set_rows(const struct ggml_compute_params * params, struct ggml_tensor * dst);

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -53,37 +53,6 @@
 
 bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor);
 
-// See https://gmplib.org/~tege/divcnst-pldi94.pdf figure 4.1.
-// Precompute mp (m' in the paper) and L such that division
-// can be computed using a multiply (high 32b of 64b result)
-// and a shift:
-//
-// n/d = (mulhi(n, mp) + n) >> L;
-struct fastdiv_vals {
-    uint32_t mp;
-    uint32_t L;
-    uint32_t d;
-    uint32_t pad;
-};
-static_assert(sizeof(fastdiv_vals) == 16, "fastdiv_vals size incorrect");
-
-static fastdiv_vals init_fastdiv_values(uint64_t d_64) {
-    GGML_ASSERT(d_64 != 0);
-    GGML_ASSERT(d_64 <= std::numeric_limits<uint32_t>::max());
-
-    uint32_t d = (uint32_t)d_64;
-
-    // compute L = ceil(log2(d));
-    uint32_t L = 0;
-    while (L < 32 && (uint32_t{ 1 } << L) < d) {
-        L++;
-    }
-
-    uint32_t mp = (uint32_t) ((uint64_t{ 1 } << 32) * ((uint64_t{ 1 } << L) - d) / d + 1);
-    // pack divisor as well to reduce error surface
-    return { mp, L, d, 0 };
-}
-
 enum GPU_FAMILY {
     ADRENO,
     INTEL,
@@ -4495,9 +4464,6 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c
             GGML_ABORT("not implemented");
     }
 
-    fastdiv_vals ne11_ = init_fastdiv_values(ne11);
-    fastdiv_vals ne12_ = init_fastdiv_values(ne12);
-
     CL_CHECK(clSetKernelArg(kernel,  0, sizeof(cl_mem),   &extra0->data_device));
     CL_CHECK(clSetKernelArg(kernel,  1, sizeof(cl_ulong), &offset0));
     CL_CHECK(clSetKernelArg(kernel,  2, sizeof(cl_mem),   &extra1->data_device));
@@ -4508,8 +4474,8 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c
     CL_CHECK(clSetKernelArg(kernel,  7, sizeof(cl_ulong), &nb01));
     CL_CHECK(clSetKernelArg(kernel,  8, sizeof(cl_ulong), &nb02));
     CL_CHECK(clSetKernelArg(kernel,  9, sizeof(cl_ulong), &nb03));
-    CL_CHECK(clSetKernelArg(kernel, 10, sizeof(fastdiv_vals), &ne11_));
-    CL_CHECK(clSetKernelArg(kernel, 11, sizeof(fastdiv_vals), &ne12_));
+    CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int),      &ne11));
+    CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int),      &ne12));
     CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb10));
     CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nb11));
     CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb12));

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

@@ -1,16 +1,5 @@
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-// v = { mp, L, d }
-inline uint fastdiv(uint n, uint4 v) {
-    uint msbs;
-    msbs = mul_hi(n, v.s0);
-    return (msbs + n) >> v.s1;
-}
-inline uint fastmod(uint n, uint4 v) {
-    uint q = fastdiv(n, v);
-    return n - q * v.s2;
-}
-
 kernel void kernel_set_rows_f32_i64(
         global char * src0,
         ulong         offset0,
@@ -22,8 +11,8 @@ kernel void kernel_set_rows_f32_i64(
         ulong         nb01,
         ulong         nb02,
         ulong         nb03,
-        uint4         ne11,
-        uint4         ne12,
+        int           ne11,
+        int           ne12,
         ulong         nb10,
         ulong         nb11,
         ulong         nb12,
@@ -44,10 +33,8 @@ kernel void kernel_set_rows_f32_i64(
         return;
     }
 
-    //int i12 = i03%ne12;
-    //int i11 = i02%ne11;
-    int i12 = fastmod(i03, ne12);
-    int i11 = fastmod(i02, ne11);
+    int i12 = i03%ne12;
+    int i11 = i02%ne11;
 
     int i10 = i01;
     long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
@@ -71,8 +58,8 @@ kernel void kernel_set_rows_f16_i64(
         ulong         nb01,
         ulong         nb02,
         ulong         nb03,
-        uint4         ne11,
-        uint4         ne12,
+        int           ne11,
+        int           ne12,
         ulong         nb10,
         ulong         nb11,
         ulong         nb12,
@@ -93,10 +80,8 @@ kernel void kernel_set_rows_f16_i64(
         return;
     }
 
-    //int i12 = i03%ne12;
-    //int i11 = i02%ne11;
-    int i12 = fastmod(i03, ne12);
-    int i11 = fastmod(i02, ne11);
+    int i12 = i03%ne12;
+    int i11 = i02%ne11;
 
     int i10 = i01;
     long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
@@ -120,8 +105,8 @@ kernel void kernel_set_rows_f32_i32(
         ulong         nb01,
         ulong         nb02,
         ulong         nb03,
-        uint4         ne11,
-        uint4         ne12,
+        int           ne11,
+        int           ne12,
         ulong         nb10,
         ulong         nb11,
         ulong         nb12,
@@ -142,10 +127,8 @@ kernel void kernel_set_rows_f32_i32(
         return;
     }
 
-    //int i12 = i03%ne12;
-    //int i11 = i02%ne11;
-    int i12 = fastmod(i03, ne12);
-    int i11 = fastmod(i02, ne11);
+    int i12 = i03%ne12;
+    int i11 = i02%ne11;
 
     int i10 = i01;
     int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
@@ -169,8 +152,8 @@ kernel void kernel_set_rows_f16_i32(
         ulong         nb01,
         ulong         nb02,
         ulong         nb03,
-        uint4         ne11,
-        uint4         ne12,
+        int           ne11,
+        int           ne12,
         ulong         nb10,
         ulong         nb11,
         ulong         nb12,
@@ -191,10 +174,8 @@ kernel void kernel_set_rows_f16_i32(
         return;
     }
 
-    //int i12 = i03%ne12;
-    //int i11 = i02%ne11;
-    int i12 = fastmod(i03, ne12);
-    int i11 = fastmod(i02, ne11);
+    int i12 = i03%ne12;
+    int i11 = i02%ne11;
 
     int i10 = i01;
     int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -6831,7 +6831,7 @@ static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context&
 
     vk_buffer d_B = d_D;
     size_t b_buf_offset = 0;
-    uint64_t b_sz = 1;
+    uint64_t b_sz = 0;
 
     if (enable_bias) {
         const ggml_tensor * add = cgraph->nodes[node_idx + 1];
@@ -6965,7 +6965,7 @@ static void ggml_vk_mul_mat_vec_p021_f16_f32(ggml_backend_vk_context * ctx, vk_c
 
     vk_buffer d_B = d_D;
     size_t b_buf_offset = 0;
-    uint64_t b_sz = 1;
+    uint64_t b_sz = 0;
 
     if (enable_bias) {
         const ggml_tensor * add = cgraph->nodes[node_idx + 1];
@@ -7101,7 +7101,7 @@ static void ggml_vk_mul_mat_vec_nc_f16_f32(ggml_backend_vk_context * ctx, vk_con
 
     vk_buffer d_B = d_D;
     size_t b_buf_offset = 0;
-    uint64_t b_sz = 1;
+    uint64_t b_sz = 0;
 
     if (enable_bias) {
         const ggml_tensor * add = cgraph->nodes[node_idx + 1];
@@ -7676,7 +7676,7 @@ static void ggml_vk_mul_mat_vec_id_q_f16(ggml_backend_vk_context * ctx, vk_conte
 
     vk_buffer d_B = d_D;
     size_t b_buf_offset = 0;
-    uint64_t b_sz = 1;
+    uint64_t b_sz = 0;
 
     if (enable_bias || enable_scale) {
         const ggml_tensor * bias = cgraph->nodes[node_idx + 1]->src[1];

ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp

@@ -18,7 +18,6 @@
 #include <algorithm>
 #include <sys/stat.h>
 #include <sys/types.h>
-#include <filesystem>
 
 #ifdef _WIN32
     #define NOMINMAX
@@ -1081,11 +1080,6 @@ int main(int argc, char** argv) {
 
     if (args.find("--glslc") != args.end()) {
         GLSLC = args["--glslc"]; // Path to glslc
-
-        if (!std::filesystem::exists(GLSLC) || !std::filesystem::is_regular_file(GLSLC)) {
-            std::cerr << "Error: glslc not found at " << GLSLC << std::endl;
-            return EXIT_FAILURE;
-        }
     }
     if (args.find("--source") != args.end()) {
         input_filepath = args["--source"]; // The shader source file to compile

src/llama-memory-recurrent.cpp

@@ -151,8 +151,7 @@ bool llama_memory_recurrent::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos
         p1 = std::numeric_limits<llama_pos>::max();
     }
 
-    // models like Mamba or RWKV can't have a state partially erased at the end
-    // of the sequence because their state isn't preserved for previous tokens
+    // models like Mamba or RWKV can't have a state partially erased
     if (seq_id >= (int64_t) size) {
         // could be fatal
         return false;
@@ -161,8 +160,8 @@ bool llama_memory_recurrent::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos
         int32_t & tail_id = cells[seq_id].tail;
         if (tail_id >= 0) {
             const auto & cell = cells[tail_id];
-            // partial intersection is invalid if it includes the final pos
-            if (0 < p0 && p0 <= cell.pos && p1 > cell.pos) {
+            // partial intersection is invalid
+            if ((0 < p0 && p0 < cell.pos) || (0 < p1 && p1 <= cell.pos)) {
                 //printf("[DEBUG] inside `llama_memory_recurrent::seq_rm`: partial intersection is invalid, so returning false\n");
                 return false;
             }

src/models/ernie4-5.cpp

@@ -1,5 +1,7 @@
 #include "models.h"
 
+
+
 llm_build_ernie4_5::llm_build_ernie4_5(const llama_model & model, const llm_graph_params & params) :
     llm_graph_context(params) {
     const int64_t n_embd_head = hparams.n_embd_head_v;
@@ -17,8 +19,6 @@ llm_build_ernie4_5::llm_build_ernie4_5(const llama_model & model, const llm_grap
 
     auto * inp_attn = build_attn_inp_kv();
 
-    ggml_tensor * inp_out_ids = build_inp_out_ids();
-
     for (int il = 0; il < n_layer; ++il) {
         ggml_tensor * inpSA = inpL;
 
@@ -67,8 +67,9 @@ llm_build_ernie4_5::llm_build_ernie4_5(const llama_model & model, const llm_grap
         }
         if (il == n_layer - 1) {
             // skip computing output for unused tokens
-            cur   = ggml_get_rows(ctx0, cur, inp_out_ids);
-            inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
+            ggml_tensor * inp_out_ids = build_inp_out_ids();
+            cur                       = ggml_get_rows(ctx0, cur, inp_out_ids);
+            inpSA                     = ggml_get_rows(ctx0, inpSA, inp_out_ids);
         }
         ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
         cb(ffn_inp, "ffn_inp", il);

src/models/openai-moe-iswa.cpp

@@ -11,8 +11,6 @@ llm_build_openai_moe_iswa::llm_build_openai_moe_iswa(const llama_model & model,
 
     auto * inp_attn = build_attn_inp_kv_iswa();
 
-    ggml_tensor * inp_out_ids = build_inp_out_ids();
-
     for (int il = 0; il < n_layer; ++il) {
         ggml_tensor * inpSA = inpL;
 
@@ -71,6 +69,7 @@ llm_build_openai_moe_iswa::llm_build_openai_moe_iswa(const llama_model & model,
         }
         if (il == n_layer - 1) {
             // skip computing output for unused tokens
+            ggml_tensor * inp_out_ids = build_inp_out_ids();
             cur   = ggml_get_rows(ctx0,   cur, inp_out_ids);
             inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
         }

tools/main/main.cpp

@@ -354,11 +354,7 @@ int main(int argc, char ** argv) {
         }
 
         // remove any "future" tokens that we might have inherited from the previous session
-        if (!llama_memory_seq_rm(mem, -1, n_matching_session_tokens, -1)) {
-            LOG_INF("%s: unable to resuse common prefix\n", __func__);
-            n_matching_session_tokens = 0;
-            llama_memory_seq_rm(mem, -1, -1, -1);
-        }
+        llama_memory_seq_rm(mem, -1, n_matching_session_tokens, -1);
     }
 
     LOG_DBG("recalculate the cached logits (check): embd_inp.size() %zu, n_matching_session_tokens %zu, embd_inp.size() %zu, session_tokens.size() %zu\n",

tools/rpc/CMakeLists.txt

@@ -2,7 +2,3 @@ set(TARGET rpc-server)
 add_executable(${TARGET} rpc-server.cpp)
 target_link_libraries(${TARGET} PRIVATE ggml)
 target_compile_features(${TARGET} PRIVATE cxx_std_17)
-
-if(LLAMA_TOOLS_INSTALL)
-    install(TARGETS ${TARGET} RUNTIME)
-endif()