metal : add ggml_metal_library_t

ggml-ci

Author: ggerganov

ggml/src/ggml-metal/CMakeLists.txt

@@ -10,7 +10,6 @@ ggml_add_backend_library(ggml-metal
                          ggml-metal-device.cpp
                          ggml-metal-common.cpp
                          ggml-metal-context.m
-                         ggml-metal-context.cpp
                          ggml-metal-ops.cpp
                         )
 

ggml/src/ggml-metal/ggml-metal-context.cpp

@@ -7,92 +7,14 @@ extern "C" {
 #endif
 
 //
-// MTLFunctionConstantValues wrapper
-//
-
-typedef struct ggml_metal_cv * ggml_metal_cv_t;
-
-ggml_metal_cv_t ggml_metal_cv_init(void);
-void ggml_metal_cv_free(ggml_metal_cv_t cv);
-
-void ggml_metal_cv_set_int32(ggml_metal_cv_t cv, int32_t value, int32_t idx);
-void ggml_metal_cv_set_bool (ggml_metal_cv_t cv, bool    value, int32_t idx);
-
-//
-// MTLComputePipelineState wrapper
-//
-
-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);
-
-void ggml_metal_pipeline_set_nsg(ggml_metal_pipeline_t pipeline, int nsg);
-int  ggml_metal_pipeline_get_nsg(ggml_metal_pipeline_t pipeline);
-
-void ggml_metal_pipeline_set_nr0(ggml_metal_pipeline_t pipeline, int nr0);
-int  ggml_metal_pipeline_get_nr0(ggml_metal_pipeline_t pipeline);
-
-void ggml_metal_pipeline_set_nr1(ggml_metal_pipeline_t pipeline, int nr1);
-int  ggml_metal_pipeline_get_nr1(ggml_metal_pipeline_t pipeline);
-
-void   ggml_metal_pipeline_set_smem(ggml_metal_pipeline_t pipeline, size_t smem);
-size_t ggml_metal_pipeline_get_smem(ggml_metal_pipeline_t pipeline);
-
-int ggml_metal_pipeline_max_theads_per_threadgroup(ggml_metal_pipeline_t pipeline);
-
-// a collection of pipelines
-typedef struct ggml_metal_pipelines * ggml_metal_pipelines_t;
-
-ggml_metal_pipelines_t ggml_metal_pipelines_init(void);
-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);
-
-//
-// backend
+// backend context
 //
 
 typedef struct ggml_metal * ggml_metal_t;
 
-ggml_metal_t ggml_metal_init(ggml_metal_device_t ctx_dev);
+ggml_metal_t ggml_metal_init(ggml_metal_device_t dev);
 void ggml_metal_free(ggml_metal_t ctx);
 
-ggml_metal_pipeline_t ggml_metal_get_pipeline    (ggml_metal_t ctx, const char * name);
-ggml_metal_pipeline_t ggml_metal_compile_pipeline(ggml_metal_t ctx, const char * base, const char * name, ggml_metal_cv_t cv);
-
 void ggml_metal_synchronize(ggml_metal_t ctx);
 
 void ggml_metal_set_tensor_async(ggml_metal_t ctx, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size);
@@ -112,28 +34,19 @@ void ggml_metal_capture_next_compute(ggml_metal_t ctx);
 
 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;
-
-    const struct ggml_metal_device_props * props_dev;
-
-    ggml_metal_encoder_t encoder;
-
-    ggml_mem_ranges_t mem_ranges;
+ggml_metal_library_t ggml_metal_graph_encoder_get_lib(ggml_metal_graph_encoder_t ctx);
+ggml_metal_encoder_t ggml_metal_graph_encoder_get_enc(ggml_metal_graph_encoder_t ctx);
+struct ggml_cgraph * ggml_metal_graph_encoder_get_gf (ggml_metal_graph_encoder_t ctx);
 
-    struct ggml_cgraph * gf;
+const struct ggml_metal_device_props * ggml_metal_graph_encoder_get_props_dev(ggml_metal_graph_encoder_t ctx);
 
-    int idx_start;
-    int idx_end;
+int ggml_metal_graph_encoder_get_idx_start(ggml_metal_graph_encoder_t ctx);
+int ggml_metal_graph_encoder_get_idx_end  (ggml_metal_graph_encoder_t ctx);
 
-    bool use_fusion;
+bool ggml_metal_graph_encoder_get_use_fusion(ggml_metal_graph_encoder_t ctx);
 
-    int debug_fusion;
-};
+int ggml_metal_graph_encoder_get_debug_fusion(ggml_metal_graph_encoder_t ctx);
+int ggml_metal_graph_encoder_get_debug_graph (ggml_metal_graph_encoder_t ctx);
 
 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);