metal : wip migrating ops

ggml-ci

Author: ggerganov

ggml/src/ggml-metal/ggml-metal-context.h

@@ -27,6 +27,9 @@ typedef struct ggml_metal_pipeline * ggml_metal_pipeline_t;
 ggml_metal_pipeline_t ggml_metal_pipeline_init(void);
 void ggml_metal_pipeline_free(ggml_metal_pipeline_t pipeline);
 
+int ggml_metal_pipeline_max_theads_per_threadgroup(ggml_metal_pipeline_t pipeline);
+
+// TODO: tmp
 void * ggml_metal_pipeline_get_obj(ggml_metal_pipeline_t pipeline);
 
 // a collection of pipelines
@@ -38,6 +41,40 @@ void ggml_metal_pipelines_free(ggml_metal_pipelines_t ppls);
 void                  ggml_metal_pipelines_add(ggml_metal_pipelines_t ppls, const char * name, ggml_metal_pipeline_t pipeline);
 ggml_metal_pipeline_t ggml_metal_pipelines_get(ggml_metal_pipelines_t ppls, const char * name);
 
+//
+// MTLCommandBuffer wrapper
+//
+
+typedef void * ggml_metal_cmd_buf_t;
+
+//
+// MTLComputeCommandEncoder wrapper
+//
+
+typedef struct ggml_metal_encoder * ggml_metal_encoder_t;
+
+ggml_metal_encoder_t ggml_metal_encoder_init(ggml_metal_cmd_buf_t cmd_buf_raw, bool concurrent);
+void ggml_metal_encoder_free(ggml_metal_encoder_t encoder);
+
+void ggml_metal_encoder_debug_group_push(ggml_metal_encoder_t encoder, const char * name);
+void ggml_metal_encoder_debug_group_pop (ggml_metal_encoder_t encoder);
+
+void ggml_metal_encoder_set_pipeline(ggml_metal_encoder_t encoder, ggml_metal_pipeline_t pipeline);
+
+void ggml_metal_encoder_set_bytes (ggml_metal_encoder_t encoder, void * data, size_t size, int idx);
+void ggml_metal_encoder_set_buffer(ggml_metal_encoder_t encoder, struct ggml_metal_buffer_id buffer, int idx);
+
+void ggml_metal_encoder_set_threadgroup_memory_size(ggml_metal_encoder_t encoder, size_t size, int idx);
+
+void ggml_metal_encoder_dispatch_threadgroups(ggml_metal_encoder_t encoder, int tg0, int tg1, int tg2, int tptg0, int tptg1, int tptg2);
+
+void ggml_metal_encoder_memory_barrier(ggml_metal_encoder_t encoder);
+
+void ggml_metal_encoder_end_encoding(ggml_metal_encoder_t encoder);
+
+// TODO: tmp
+void * ggml_metal_encoder_get_obj(ggml_metal_encoder_t encoder);
+
 //
 // backend
 //
@@ -63,6 +100,341 @@ void ggml_metal_set_abort_callback  (ggml_metal_t ctx, ggml_abort_callback abort
 bool ggml_metal_supports_family     (ggml_metal_t ctx, int family);
 void ggml_metal_capture_next_compute(ggml_metal_t ctx);
 
+//
+// graph encoder
+//
+
+typedef struct ggml_metal_graph_encoder * ggml_metal_graph_encoder_t;
+
+// TODO: tmp
+#include "ggml-metal-common.h"
+
+// TODO: tmp
+struct ggml_metal_graph_encoder {
+    ggml_metal_t ctx;
+
+    ggml_metal_encoder_t encoder;
+
+    ggml_mem_ranges_t mem_ranges;
+
+    struct ggml_cgraph * gf;
+
+    int idx_start;
+    int idx_end;
+
+    bool use_fusion;
+
+    int debug_fusion;
+};
+
+bool ggml_metal_graph_encoder_concurrency_reset(ggml_metal_graph_encoder_t ctx);
+bool ggml_metal_graph_encoder_concurrency_check(ggml_metal_graph_encoder_t ctx, const struct ggml_tensor * node);
+bool ggml_metal_graph_encoder_concurrency_add  (ggml_metal_graph_encoder_t ctx, const struct ggml_tensor * node);
+
+// TODO: tmp
+enum ggml_metal_pipeline_type {
+    GGML_METAL_PIPELINE_TYPE_ADD_ID,
+    GGML_METAL_PIPELINE_TYPE_REPEAT_F32,
+    GGML_METAL_PIPELINE_TYPE_REPEAT_F16,
+    GGML_METAL_PIPELINE_TYPE_REPEAT_I32,
+    GGML_METAL_PIPELINE_TYPE_REPEAT_I16,
+    GGML_METAL_PIPELINE_TYPE_SCALE,
+    GGML_METAL_PIPELINE_TYPE_SCALE_4,
+    GGML_METAL_PIPELINE_TYPE_CLAMP,
+    GGML_METAL_PIPELINE_TYPE_TANH,
+    GGML_METAL_PIPELINE_TYPE_RELU,
+    GGML_METAL_PIPELINE_TYPE_SIGMOID,
+    GGML_METAL_PIPELINE_TYPE_GELU,
+    GGML_METAL_PIPELINE_TYPE_GELU_4,
+    GGML_METAL_PIPELINE_TYPE_GELU_ERF,
+    GGML_METAL_PIPELINE_TYPE_GELU_ERF_4,
+    GGML_METAL_PIPELINE_TYPE_GELU_QUICK,
+    GGML_METAL_PIPELINE_TYPE_GELU_QUICK_4,
+    GGML_METAL_PIPELINE_TYPE_SILU,
+    GGML_METAL_PIPELINE_TYPE_SILU_4,
+    GGML_METAL_PIPELINE_TYPE_ELU,
+    GGML_METAL_PIPELINE_TYPE_ABS,
+    GGML_METAL_PIPELINE_TYPE_SGN,
+    GGML_METAL_PIPELINE_TYPE_STEP,
+    GGML_METAL_PIPELINE_TYPE_HARDSWISH,
+    GGML_METAL_PIPELINE_TYPE_HARDSIGMOID,
+    GGML_METAL_PIPELINE_TYPE_EXP,
+    GGML_METAL_PIPELINE_TYPE_SOFT_MAX_F16,
+    GGML_METAL_PIPELINE_TYPE_SOFT_MAX_F16_4,
+    GGML_METAL_PIPELINE_TYPE_SOFT_MAX_F32,
+    GGML_METAL_PIPELINE_TYPE_SOFT_MAX_F32_4,
+    GGML_METAL_PIPELINE_TYPE_DIAG_MASK_INF,
+    GGML_METAL_PIPELINE_TYPE_DIAG_MASK_INF_8,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_F32,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_F16,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_BF16,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q4_0,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q4_1,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q5_0,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q5_1,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q8_0,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_MXFP4,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q2_K,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q3_K,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q4_K,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q5_K,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_Q6_K,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ2_XXS,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ2_XS,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ3_XXS,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ3_S,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ2_S,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ1_S,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ1_M,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ4_NL,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_IQ4_XS,
+    GGML_METAL_PIPELINE_TYPE_GET_ROWS_I32,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_F32,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_F16,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_BF16,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_Q8_0,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_Q4_0,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_Q4_1,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_Q5_0,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_Q5_1,
+    GGML_METAL_PIPELINE_TYPE_SET_ROWS_IQ4_NL,
+    GGML_METAL_PIPELINE_TYPE_L2_NORM,
+    GGML_METAL_PIPELINE_TYPE_GROUP_NORM,
+    GGML_METAL_PIPELINE_TYPE_NORM,
+    GGML_METAL_PIPELINE_TYPE_SSM_CONV_F32,
+    GGML_METAL_PIPELINE_TYPE_SSM_SCAN_F32,
+    GGML_METAL_PIPELINE_TYPE_SSM_SCAN_F32_GROUP,
+    GGML_METAL_PIPELINE_TYPE_RWKV_WKV6_F32,
+    GGML_METAL_PIPELINE_TYPE_RWKV_WKV7_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F32_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F32_F32_C4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F16_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F16_F32_C4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F16_F32_1ROW,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F16_F32_L4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_F16_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_BF16_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_BF16_F32_C4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_BF16_F32_1ROW,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_BF16_F32_L4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_BF16_BF16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q4_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q4_1_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q5_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q5_1_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q8_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_MXFP4_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F32_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F32_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F32_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F32_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F16_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F16_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F16_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_F16_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_0_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_0_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_0_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_0_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_1_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_1_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_1_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_1_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_0_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_0_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_0_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_0_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_1_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_1_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_1_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_1_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q8_0_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q8_0_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q8_0_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q8_0_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_MXFP4_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_MXFP4_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_MXFP4_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_MXFP4_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_K_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_K_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_K_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q4_K_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_K_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_K_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_K_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q5_K_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q6_K_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q6_K_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q6_K_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_Q6_K_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_IQ4_NL_F32_R1_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_IQ4_NL_F32_R1_3,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_IQ4_NL_F32_R1_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_EXT_IQ4_NL_F32_R1_5,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q2_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q3_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q4_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q5_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_Q6_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ2_XXS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ2_XS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ3_XXS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ3_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ2_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ1_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ1_M_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ4_NL_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_IQ4_XS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_F32_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_F16_F32,
+  //GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_F16_F32_1ROW,
+  //GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_F16_F32_L4,
+  //GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_F16_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_BF16_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q4_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q4_1_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q5_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q5_1_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q8_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_MXFP4_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q2_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q3_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q4_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q5_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_Q6_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ2_XXS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ2_XS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ3_XXS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ3_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ2_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ1_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ1_M_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ4_NL_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MV_ID_IQ4_XS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_F32_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_F16_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_BF16_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q4_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q4_1_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q5_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q5_1_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q8_0_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_MXFP4_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q2_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q3_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q4_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q5_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_Q6_K_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ2_XXS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ2_XS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ3_XXS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ3_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ2_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ1_S_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ1_M_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ4_NL_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_IQ4_XS_F32,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_1,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_2,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_4,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_6,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_8,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_10,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MAP0_F16_NE20_16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_F32_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_F16_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_BF16_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q4_0_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q4_1_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q5_0_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q5_1_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q8_0_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_MXFP4_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q2_K_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q3_K_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q4_K_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q5_K_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_Q6_K_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ2_XXS_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ2_XS_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ3_XXS_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ3_S_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ2_S_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ1_S_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ1_M_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ4_NL_F16,
+    GGML_METAL_PIPELINE_TYPE_MUL_MM_ID_IQ4_XS_F16,
+    GGML_METAL_PIPELINE_TYPE_ROPE_NORM_F32,
+    GGML_METAL_PIPELINE_TYPE_ROPE_NORM_F16,
+    GGML_METAL_PIPELINE_TYPE_ROPE_MULTI_F32,
+    GGML_METAL_PIPELINE_TYPE_ROPE_MULTI_F16,
+    GGML_METAL_PIPELINE_TYPE_ROPE_VISION_F32,
+    GGML_METAL_PIPELINE_TYPE_ROPE_VISION_F16,
+    GGML_METAL_PIPELINE_TYPE_ROPE_NEOX_F32,
+    GGML_METAL_PIPELINE_TYPE_ROPE_NEOX_F16,
+    GGML_METAL_PIPELINE_TYPE_IM2COL_F16,
+    GGML_METAL_PIPELINE_TYPE_IM2COL_F32,
+    GGML_METAL_PIPELINE_TYPE_IM2COL_EXT_F16,
+    GGML_METAL_PIPELINE_TYPE_IM2COL_EXT_F32,
+    GGML_METAL_PIPELINE_TYPE_CONV_TRANSPOSE_1D_F32_F32,
+    GGML_METAL_PIPELINE_TYPE_CONV_TRANSPOSE_1D_F16_F32,
+    GGML_METAL_PIPELINE_TYPE_UPSCALE_F32,
+    GGML_METAL_PIPELINE_TYPE_PAD_F32,
+    GGML_METAL_PIPELINE_TYPE_PAD_REFLECT_1D_F32,
+    GGML_METAL_PIPELINE_TYPE_ARANGE_F32,
+    GGML_METAL_PIPELINE_TYPE_TIMESTEP_EMBEDDING_F32,
+    GGML_METAL_PIPELINE_TYPE_ARGSORT_F32_I32_ASC,
+    GGML_METAL_PIPELINE_TYPE_ARGSORT_F32_I32_DESC,
+    GGML_METAL_PIPELINE_TYPE_LEAKY_RELU_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_F16,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_BF16,
+    GGML_METAL_PIPELINE_TYPE_CPY_F16_F16,
+    GGML_METAL_PIPELINE_TYPE_CPY_F16_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_BF16_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_BF16_BF16,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_I32,
+    GGML_METAL_PIPELINE_TYPE_CPY_I32_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_Q8_0,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_Q4_0,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_Q4_1,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_Q5_0,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_Q5_1,
+    GGML_METAL_PIPELINE_TYPE_CPY_F32_IQ4_NL,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q4_0_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q4_0_F16,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q4_1_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q4_1_F16,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q5_0_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q5_0_F16,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q5_1_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q5_1_F16,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q8_0_F32,
+    GGML_METAL_PIPELINE_TYPE_CPY_Q8_0_F16,
+    GGML_METAL_PIPELINE_TYPE_CONCAT,
+    GGML_METAL_PIPELINE_TYPE_SQR,
+    GGML_METAL_PIPELINE_TYPE_SQRT,
+    GGML_METAL_PIPELINE_TYPE_SIN,
+    GGML_METAL_PIPELINE_TYPE_COS,
+    GGML_METAL_PIPELINE_TYPE_NEG,
+    GGML_METAL_PIPELINE_TYPE_REGLU,
+    GGML_METAL_PIPELINE_TYPE_GEGLU,
+    GGML_METAL_PIPELINE_TYPE_SWIGLU,
+    GGML_METAL_PIPELINE_TYPE_SWIGLU_OAI,
+    GGML_METAL_PIPELINE_TYPE_GEGLU_ERF,
+    GGML_METAL_PIPELINE_TYPE_GEGLU_QUICK,
+    GGML_METAL_PIPELINE_TYPE_SUM_ROWS,
+    GGML_METAL_PIPELINE_TYPE_MEAN,
+    GGML_METAL_PIPELINE_TYPE_POOL_2D_AVG_F32,
+    GGML_METAL_PIPELINE_TYPE_POOL_2D_MAX_F32,
+    GGML_METAL_PIPELINE_TYPE_ARGMAX,
+
+    GGML_METAL_PIPELINE_TYPE_COUNT
+};
+
+ggml_metal_pipeline_t ggml_metal_get_pipeline_by_type(ggml_metal_t ctx, enum ggml_metal_pipeline_type type);
+
 #ifdef __cplusplus
 }
 #endif