Repository: ggml-org/llama.cpp
Pull request: sycl: GGML_SYCL_DISABLE_OPT on by default for all Intel Devices

This PR proposes moving the env variable `GGML_SYCL_DISABLE_OPT` from default ON (reorder disabled) to default OFF (reorders enabled). This would allow easier testing on newer hardware, without needing to modify the list of devices which support the reorder feature.
Concerns regarding the reorder feature from #13254 have been resolved.
Regarding performance on older devices an amendment has been made to the README to suggest disabling the feature.

Below are performance's runs on 2 different models showing performance improvements of having the feature enabled:
## Llama2-7B Q4_0 PVC

With reorder

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           pp512 |      2618.08 ± 13.71 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           tg128 |         74.49 ± 0.25 |

```

Without

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           pp512 |      2611.12 ± 25.25 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           tg128 |         36.02 ± 0.07 |

```

## gemma2 2B Q4_K PVC

With Reorder

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           pp512 |      7765.55 ± 60.15 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           tg128 |         99.06 ± 0.12 |
```

Without

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           pp512 |     7766.70 ± 145.88 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           tg128 |         89.94 ± 0.28 |
```

## Llama2-7B Q4_0 ARC-A770

With reorder

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           pp512 |       1711.58 ± 3.83 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           tg128 |         34.16 ± 0.22 |

```

Without

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           pp512 |       1711.60 ± 1.30 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           tg128 |         29.97 ± 0.22 |

```

## gemma2 2B Q4_K ARC-A770

With Reorder

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           pp512 |       3645.32 ± 4.61 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           tg128 |         38.23 ± 0.13 |

```

Without

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           pp512 |       3639.22 ± 7.44 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           tg128 |         35.55 ± 0.14 |

```

## Llama2-7B Q4_0 Lunar Lake

With reorder

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           pp512 |       258.61 ± 24.36 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           tg128 |         19.85 ± 0.04 |

```

Without

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           pp512 |        470.37 ± 1.56 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |       8 |           tg128 |         12.92 ± 0.99 |

```

## gemma2 2B Q4_K Lunar Lake

With Reorder

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           pp512 |       613.06 ± 21.23 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           tg128 |         29.01 ± 0.32 |

```

Without

```cpp
| model                          |       size |     params | backend    | ngl | threads |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           pp512 |       643.79 ± 87.98 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       8 |           tg128 |         24.78 ± 0.27 |

```

## Llama2-7B Q4_0 Intel B580

With reorder

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           pp512 |      2162.42 ± 13.20 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           tg128 |         66.66 ± 0.21 |
```

Without

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           pp512 |       2168.30 ± 6.27 |
| llama 7B Q4_0                  |   3.56 GiB |     6.74 B | SYCL       |  99 |           tg128 |         39.67 ± 0.09 |

```

## gemma2 2B Q4_K Intel B580

With Reorder

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           pp512 |      5685.01 ± 21.13 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           tg128 |         87.73 ± 1.75 |
```

Without

```cpp
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           pp512 |      5678.08 ± 18.61 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |           tg128 |         66.92 ± 0.62 |
```


Changed files:
filename: docs/backend/SYCL.md
patch: @@ -757,7 +757,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512
 | Name              | Value            | Function                                                                                                                  |
 |-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------|
 | GGML_SYCL_DEBUG   | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG                                                                             |
-| GGML_SYCL_DISABLE_OPT | 0 (default) or 1 | Disable optimize features based on Intel GPU type, to compare the performance increase |
+| GGML_SYCL_DISABLE_OPT | 0 (default) or 1 | Disable optimize features for Intel GPUs. (Recommended to 1 for intel devices older than Gen 10) |
 | GGML_SYCL_DISABLE_GRAPH | 0 or 1 (default) | Disable running computations through SYCL Graphs feature. Disabled by default because graph performance isn't yet better than non-graph performance. |
 | GGML_SYCL_DISABLE_DNN | 0 (default) or 1 | Disable running computations through oneDNN and always use oneMKL. |
 | ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.<br>Recommended to use when --split-mode = layer |

filename: ggml/src/ggml-sycl/common.hpp
patch: @@ -199,7 +199,7 @@ struct sycl_device_info {
     // size_t  smpb;               // max. shared memory per block
     bool    vmm;                // virtual memory support
     size_t  total_vram;
-    sycl_hw_info hw_info;
+    //sycl_hw_info hw_info;     \\ device id and aarch, currently not used
     optimize_feature opt_feature;
 };
 
@@ -286,29 +286,6 @@ struct ggml_tensor_extra_gpu {
 
 void release_extra_gpu(ggml_tensor_extra_gpu * extra, std::vector<queue_ptr> streams={});
 
-inline optimize_feature check_gpu_optimize_feature(syclex::architecture &arch) {
-    optimize_feature opt;
-
-    opt.reorder =
-        (arch == syclex::architecture::intel_gpu_dg1 ||
-         arch == syclex::architecture::intel_gpu_acm_g10 ||
-         arch == syclex::architecture::intel_gpu_acm_g11 ||
-         arch == syclex::architecture::intel_gpu_acm_g12 ||
-         arch == syclex::architecture::intel_gpu_pvc ||
-         arch == syclex::architecture::intel_gpu_pvc_vg ||
-         arch == syclex::architecture::intel_gpu_mtl_u ||
-         arch == syclex::architecture::intel_gpu_mtl_s ||
-         arch == syclex::architecture::intel_gpu_mtl_h ||
-         arch == syclex::architecture::intel_gpu_arl_u ||
-         arch == syclex::architecture::intel_gpu_arl_s ||
-         arch == syclex::architecture::intel_gpu_arl_h ||
-         arch == syclex::architecture::intel_gpu_bmg_g21 ||
-         arch == syclex::architecture::intel_gpu_lnl_m
-        );
-
-    return opt;
-}
-
 namespace sycl_ex = sycl::ext::oneapi::experimental;
 struct ggml_backend_sycl_context {
     int device;

filename: ggml/src/ggml-sycl/ggml-sycl.cpp
patch: @@ -83,9 +83,7 @@ static ggml_sycl_device_info ggml_sycl_init() {
 
         info.devices[i].cc =
             100 * prop.get_major_version() + 10 * prop.get_minor_version();
-        info.devices[i].hw_info = get_device_hw_info(&device);
-        info.devices[i].opt_feature = check_gpu_optimize_feature(info.devices[i].hw_info.arch);
-
+        info.devices[i].opt_feature.reorder = !device.ext_oneapi_architecture_is(syclex::arch_category::intel_gpu);
         info.max_work_group_sizes[i] = prop.get_max_work_group_size();
     }
 
@@ -195,7 +193,7 @@ static void ggml_check_sycl() try {
 
     if (!initialized) {
         g_ggml_sycl_debug = get_sycl_env("GGML_SYCL_DEBUG", 0);
-        g_ggml_sycl_disable_optimize= get_sycl_env("GGML_SYCL_DISABLE_OPT", 1);
+        g_ggml_sycl_disable_optimize = get_sycl_env("GGML_SYCL_DISABLE_OPT", 0);
         g_ggml_sycl_disable_graph = get_sycl_env("GGML_SYCL_DISABLE_GRAPH", 1);
         g_ggml_sycl_disable_dnn = get_sycl_env("GGML_SYCL_DISABLE_DNN", 0);
         g_ggml_sycl_prioritize_dmmv = get_sycl_env("GGML_SYCL_PRIORITIZE_DMMV", 0);

filename: ggml/src/ggml-sycl/sycl_hw.cpp
patch: @@ -1,6 +1,7 @@
 #include "sycl_hw.hpp"
 
-
+// TODO: currently not used
+/*
 sycl_hw_info get_device_hw_info(sycl::device *device_ptr) {
   sycl_hw_info res;
   int32_t id = device_ptr->get_info<sycl::ext::intel::info::device::device_id>();
@@ -11,3 +12,4 @@ sycl_hw_info get_device_hw_info(sycl::device *device_ptr) {
 
   return res;
 }
+*/

filename: ggml/src/ggml-sycl/sycl_hw.hpp
patch: @@ -10,6 +10,8 @@
 
 namespace syclex = sycl::ext::oneapi::experimental;
 
+// TODO: currently not used
+/*
 struct sycl_hw_info {
   syclex::architecture arch;
   int32_t device_id;
@@ -18,6 +20,7 @@ struct sycl_hw_info {
 bool is_in_vector(std::vector<int> &vec, int item);
 
 sycl_hw_info get_device_hw_info(sycl::device *device_ptr);
+*/
 
 
 #endif // SYCL_HW_HPP