bulk refactoring task profile and related to run CL GPU offloading.

* removed ggml_task_backend, infavour of ggml_task_profile.runner and newly added id and name.
* extracted mul_mat blas codes into ggml_compute_forward_mul_mat_blas,
  thus align with CUDA/CL a bit more and make it easier to fix profile and run tune.
* rewrote task profile and update/add some cuda/cl codes, finnaly made CL GPU offloading work.
* misc minor fix/update to tune, the data format was changed.

Author: mqy

examples/mulmat-tune/README.md

@@ -214,26 +214,19 @@ The following results are generated with Accelerate compiled.
 **Example**
 
 ```
-5 3B 2 6 1
-
-3200 3200  2 0 3 10
-16 0 0 0  16 1 0 1   0 0 0 0
-16 1 0 2  17 0 1 0   0 0 0 0
- 0 0 0 0  34 0 1 0   0 0 0 0
-   1        1      793 0     9103     2102 0 0     6014 0
-   2        2     1591 0     8034     2305 0 0    30982 0
-   4        4     2236 0     6476     2484 0 0    31388 0
-   8        7     4161 0     6623     2389 0 0    29204 0
-  16       15     8339 0     6434     2752 0 0    34303 0
-  32       32    16919 0     6915     3651 0 0    42511 0
-  64      200    34270 0     6574     4528 0 0    68212 0
- 128      188    69400 0     6325     6839 0 0    74437 0
- 256      303   134597 0     6168    11544 0 0   110180 0
- 512      687   279685 0     6337    29712 0 0   159728 0
-
-3200 8640  2 0 2 10
-
- ...
+[tune] done, elapsed time: 0 seconds.
+10 xB 12 4 2
+
+1024 1024 12 0 2 4
+100 110 000 1 CPU
+110 101 000 2 BLAS
+   1       11      309 0     1234       90 0
+   2       23      654 0     1359      215 0
+   4       44     1283 0     1362      421 0
+   8       85     2341 0     1357      347 0
+
+1024 2048 12 0 2 4
+...
 
  ```
 
@@ -249,17 +242,17 @@ shape+
 # head
 version: 1
 model: "3B" | "7B" | "13B" | "30B" | "65B"
-ggml_ftype: 0 - 4, 7 - 14
+ggml_ftype: 0 - 3, 7 - 14
 n_shapes: number of shapes
 n_threads: number of threads
 
-shape := N K  m_num n_profiles
-task_conf_profile+
+shape := N K  src0_ggml_type src1_ggml_type n_profiles m_num
+task_profile+
 bench_item+
 
-task_conf_profile: stage_conf(init) stage_conf(compute) stage_conf(finalize)
-stage_conf: backend parallel wait
-backend: 0 (NONE) | 16 (CPU) | 17 (CPU_BLAS) | 32 (GPU) | 33 (GPU_CUDA) | 34 (GPU_CL)
+task_profile: stage_conf(init) stage_conf(compute) stage_conf(finalize) id name
+stage_conf(bitmap): valid parallel wait
+valid: 0 (false) | 1 (true)
 parallel: 0 (false) | 1 (true)
 wait: 0 (false) | 1 (true)
 

examples/mulmat-tune/mulmat-tune.cpp

@@ -111,6 +111,11 @@ static void usage(char *prog) {
 }
 
 int main(int argc, char **argv) {
+    if (!ggml_cpu_has_blas()) {
+        fprintf(stderr, "error: this program is not built with BLAS.\n");
+        return 1;
+    }
+
     if (argc == 2) {
         if (strcmp(argv[1], "-h") == 0 || strcmp(argv[1], "--help") == 0) {
             usage(argv[0]);

ggml-cuda.cu

@@ -2207,17 +2207,12 @@ void ggml_cuda_rms_norm(const ggml_tensor * src0, const ggml_tensor * src1, ggml
     ggml_cuda_op(src0, src1, dst, ggml_cuda_op_rms_norm, true, true);
 }
 
-bool ggml_cuda_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst) {
-    const int64_t ne10 = src1->ne[0];
-
-    const int64_t ne0 = dst->ne[0];
-    const int64_t ne1 = dst->ne[1];
-
+// NOTE: don't check matrix size, otherwise mul_mat tune will fail to run.
+static bool ggml_cuda_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst) {
     // TODO: find the optimal values for these
     if ((src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) &&
         src1->type == GGML_TYPE_F32 &&
-        dst->type == GGML_TYPE_F32 &&
-        (ne0 >= 32 && ne1 >= 32 && ne10 >= 32)) {
+        dst->type == GGML_TYPE_F32) {
         return true;
     }
 
@@ -2539,11 +2534,17 @@ void ggml_cuda_free_scratch() {
     g_scratch_buffer = nullptr;
 }
 
-bool ggml_cuda_compute_forward(struct ggml_compute_params * params, struct ggml_tensor * tensor){
-    ggml_cuda_func_t func;
-    const bool any_on_device = tensor->backend == GGML_BACKEND_GPU
+bool ggml_cuda_is_gpu_offloading(struct ggml_tensor * tensor) {
+    GGML_ASSERT(tensor);
+    GGML_ASSERT(tensor->src0);
+    return tensor->backend == GGML_BACKEND_GPU
         || tensor->src0->backend == GGML_BACKEND_GPU || tensor->src0->backend == GGML_BACKEND_GPU_SPLIT
         || (tensor->src1 != nullptr && tensor->src1->backend == GGML_BACKEND_GPU);
+}
+
+bool ggml_cuda_compute_forward(struct ggml_compute_params * params, struct ggml_tensor * tensor){
+    ggml_cuda_func_t func;
+    const bool any_on_device = is_gpu_offloading(tensor);
 
     switch (tensor->op) {
         case GGML_OP_ADD:
@@ -2571,7 +2572,7 @@ bool ggml_cuda_compute_forward(struct ggml_compute_params * params, struct ggml_
             func = ggml_cuda_rms_norm;
             break;
         case GGML_OP_MUL_MAT:
-            if (!any_on_device/* && !ggml_cuda_can_mul_mat(tensor->src0, tensor->src1, tensor)*/) {
+            if (!any_on_device && !ggml_cuda_can_mul_mat(tensor->src0, tensor->src1, tensor)) {
                 return false;
             }
             func = ggml_cuda_mul_mat;

ggml-cuda.h

@@ -16,7 +16,7 @@ void   ggml_init_cublas(void);
 void   ggml_cuda_set_tensor_split(const float * tensor_split);
 
 void   ggml_cuda_mul(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst);
-bool   ggml_cuda_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst);
+bool   ggml_cuda_is_gpu_offloading(const struct ggml_tensor * src0);
 size_t ggml_cuda_mul_mat_get_wsize(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst);
 void   ggml_cuda_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst, void * wdata, size_t wsize);
 

ggml-opencl.cpp

@@ -1589,18 +1589,17 @@ static void ggml_cl_mul_mat_q_f32(const ggml_tensor * src0, const ggml_tensor *
     }
 }
 
+bool ggml_cl_is_gpu_offloading(struct ggml_tensor * tensor) {
+    GGML_ASSERT(tensor);
+    return (tensor->src0 && tensor->src0->backend == GGML_BACKEND_GPU) ||
+        (tensor->src1 && tensor->src1->backend == GGML_BACKEND_GPU);
+}
 
-bool ggml_cl_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst) {
-    const int64_t ne10 = src1->ne[0];
-
-    const int64_t ne0 = dst->ne[0];
-    const int64_t ne1 = dst->ne[1];
-
-    // TODO: find the optimal values for these
+// NOTE: don't check matrix size, otherwise mul_mat tune will fail to run.
+static bool ggml_cl_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst) {
     if ((src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) &&
         src1->type == GGML_TYPE_F32 &&
-        dst->type == GGML_TYPE_F32 /*&&
-        ((ne0 >= 32 && ne1 >= 32 && ne10 >= 32) || src0->backend == GGML_BACKEND_GPU)*/) {
+        dst->type == GGML_TYPE_F32) {
         return true;
     }
 

ggml-opencl.h

@@ -9,7 +9,7 @@ extern "C" {
 void ggml_cl_init(void);
 
 void   ggml_cl_mul(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst);
-bool   ggml_cl_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst);
+bool   ggml_cl_is_gpu_offloading(struct ggml_tensor * tensor);
 size_t ggml_cl_mul_mat_get_wsize(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst);
 void   ggml_cl_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst, void * wdata, size_t wsize);
 

ggml-threading.c

@@ -376,7 +376,7 @@ ggml_thread_ret_t ggml_threading_graph_compute_thread(void *data) {
 
     struct ggml_compute_state_shared *shared = state->shared;
     GGML_ASSERT(shared);
-    GGML_ASSERT(shared->task_runner);
+    //GGML_ASSERT(shared->task_runner);
 
     shared->n_ready++;
 
@@ -397,7 +397,7 @@ ggml_thread_ret_t ggml_threading_graph_compute_thread(void *data) {
                                            : shared->task_runner;
             enum ggml_compute_error err = runner(&state->params, state->node);
 
-            GGML_ASSERT(err == GGML_COMPUTE_OK);
+            GGML_ASSERT(err == GGML_COMPUTE_OK || err == GGML_COMPUTE_FALLBACK);
 
             ggml_spin_lock(&shared->spin);
 
@@ -430,7 +430,7 @@ ggml_threading_compute_tensor(struct ggml_threading_context *ctx,
                               size_t wsize) {
     GGML_ASSERT(ctx);
     GGML_ASSERT(node);
-    GGML_ASSERT(ctx->shared.task_runner);
+    // GGML_ASSERT(ctx->shared.task_runner);
 
     ggml_task_runner *runner = ctx->shared.task_runner;
     if (node->task_profile.runner) {
@@ -448,7 +448,7 @@ ggml_threading_compute_tensor(struct ggml_threading_context *ctx,
     memset(&params, 0, sizeof(struct ggml_compute_params));
 
     for (int type = GGML_TASK_INIT; type <= GGML_TASK_FINALIZE; type++) {
-        if (node->task_profile.stages[type].backend == GGML_TASK_BACKEND_NONE) {
+        if (!node->task_profile.stages[type].valid) {
             continue;
         }
 
@@ -519,18 +519,17 @@ ggml_threading_compute_tensor(struct ggml_threading_context *ctx,
             if (err == GGML_COMPUTE_FALLBACK) {
                 PRINT_DEBUG("[main] fallback from profile, id=%d\n",
                             node->task_profile.id);
-                GGML_ASSERT(node->task_profile.stages[1].backend >
-                            GGML_TASK_BACKEND_CPU);
+                GGML_ASSERT(node->task_profile.id > 1);
 
                 struct ggml_task_profile profiles[GGML_MAX_TASK_PROFILES];
                 int n = ggml_get_task_profiles(node, profiles);
                 GGML_ASSERT(n > 0);
-                GGML_ASSERT(profiles[0].stages[1].backend ==
-                            GGML_TASK_BACKEND_CPU);
+                GGML_ASSERT(profiles[0].id == 1);
 
                 memcpy(&node->task_profile, &profiles[0],
-                       sizeof(struct ggml_task_profile));
+                    sizeof(struct ggml_task_profile));
                 runner = ctx->shared.task_runner;
+                GGML_ASSERT(runner);
 
                 goto START;
             }

ggml-threading.h

@@ -29,7 +29,9 @@ typedef ggml_thread_ret_t(ggml_threading_thread_runner)(void *data);
 // thread: optional OS thread runner, default value:
 // `ggml_threading_graph_compute_thread`.
 //
-// features: optional for configure
+// task_runner: default task runner, nullable wheen tensor.runner is not NULL.
+//              Overridden by tensor.runner.
+// features: configure threading behaviour, optional.
 // threading additional features. see `ggml_threading_feature`, default 0.
 //
 // stages_time: optional for collecting per-stage wall clock time.
@@ -51,12 +53,6 @@ enum ggml_compute_error
 ggml_threading_compute_tensor(struct ggml_threading_context *ctx,
                               struct ggml_tensor *node, void *wdata,
                               size_t wsize);
-
-// This is an experimental functionality for mulmat tune, as a thin wrapper.
-enum ggml_compute_error
-ggml_compute_forward_wrapper(const struct ggml_compute_params *params,
-                             struct ggml_tensor *tensor);
-
 #ifdef __cplusplus
 }
 #endif