Repository: ggml-org/llama.cpp
Pull request: sycl : Overcoming workaround for mmap() allocation on Windows

This PR removes the usage of a workaround for mmap bug on some Intel GPUs on Linux. The bug is not present on Windows, so there is no meaning of having it in place.
This causes a small split in the codebase according to the OS in use, but it shows good performance improvements.

The work introduced here is based on #13109

**N.B All numbers assessed with `GGML_SYCL_DISABLE_OPT=0`**

#### Lunar Lake's performance (this PR)

| model                          |       size |     params | backend    | ngl |    sm |          test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ----: | ------------: | -------------------: |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |  none |         pp512 |       1330.42 ± 6.59 |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |  none |         tg128 |         58.92 ± 0.46 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |  none |         pp512 |      2044.01 ± 13.08 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |  none |         tg128 |         44.47 ± 0.13 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |  none |         pp512 |        320.23 ± 0.97 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |  none |         tg128 |         22.66 ± 0.02 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |  none |         pp512 |        533.16 ± 1.41 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |  none |         tg128 |         15.41 ± 0.44 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |  none |         pp512 |       1402.31 ± 7.56 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |  none |         tg128 |         28.55 ± 0.06 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |  none |         pp512 |        502.78 ± 1.02 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |  none |         tg128 |         35.83 ± 0.07 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |  none |         pp512 |        807.02 ± 2.71 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |  none |         tg128 |         23.57 ± 0.08 |

build: 0e1009f16 (5334)

#### Lunar Lake's performance (#13109)

| model                          |       size |     params | backend    | ngl |    sm | mmap |          test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ----: | ---: | ------------: | -------------------: |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |  none |    0 |         pp512 |       1323.21 ± 8.43 |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |  none |    0 |         tg128 |         52.47 ± 0.42 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |  none |    0 |         pp512 |       1994.78 ± 6.69 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |  none |    0 |         tg128 |         40.50 ± 0.10 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |  none |    0 |         pp512 |        297.47 ± 0.49 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |  none |    0 |         tg128 |         21.58 ± 0.08 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |  none |    0 |         pp512 |        499.53 ± 2.32 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |  none |    0 |         tg128 |         15.54 ± 0.31 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |  none |    0 |         pp512 |        907.84 ± 0.56 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |  none |    0 |         tg128 |         27.54 ± 0.09 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |  none |    0 |         pp512 |        477.35 ± 0.33 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |  none |    0 |         tg128 |         33.95 ± 0.07 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |  none |    0 |         pp512 |        757.61 ± 1.53 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |  none |    0 |         tg128 |         21.80 ± 0.32 |

build: f7e7d2a5d (5331)

#### Battlemage(B580) performance (this PR)

| model                          |       size |     params | backend    | ngl | threads |    sm |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | ----: | --------------: | -------------------: |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |       5 |  none |           pp512 |      7314.80 ± 23.23 |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |       5 |  none |           tg128 |         71.10 ± 2.21 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |       5 |  none |           pp512 |      7419.09 ± 27.47 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |       5 |  none |           tg128 |         88.57 ± 0.12 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |       5 |  none |           pp512 |       2147.78 ± 6.70 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |       5 |  none |           tg128 |         40.59 ± 0.07 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |       5 |  none |           pp512 |       2189.34 ± 2.19 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |       5 |  none |           tg128 |         38.32 ± 0.02 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       5 |  none |           pp512 |      5605.63 ± 22.70 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       5 |  none |           tg128 |         72.54 ± 0.29 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |       5 |  none |           pp512 |       3002.45 ± 4.25 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |       5 |  none |           tg128 |         62.49 ± 0.04 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |       5 |  none |           pp512 |       3103.20 ± 3.79 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |       5 |  none |           tg128 |         58.64 ± 0.01 |

build: 0e1009f1 (5334)

#### Battlemage(B580) performance(#13109 )

| model                          |       size |     params | backend    | ngl | threads |    sm | mmap |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ------: | ----: | ---: | --------------: | -------------------: |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |      7067.24 ± 53.67 |
| qwen2 1.5B Q4_0                | 1013.62 MiB |     1.78 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         64.51 ± 0.33 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |      7132.89 ± 28.96 |
| qwen2 1.5B Q4_K - Medium       |   1.04 GiB |     1.78 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         78.58 ± 0.19 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |       2109.49 ± 2.46 |
| llama 7B Q4_0                  |   3.57 GiB |     6.74 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         38.37 ± 0.11 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |       2143.62 ± 0.99 |
| llama 7B Q4_K - Medium         |   3.80 GiB |     6.74 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         36.33 ± 0.03 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |      5322.20 ± 22.77 |
| gemma2 2B Q4_K - Medium        |   1.59 GiB |     2.61 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         64.48 ± 0.08 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |       2936.43 ± 7.73 |
| phi3 3B Q4_0                   |   2.03 GiB |     3.82 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         57.50 ± 0.11 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |       5 |  none |    0 |           pp512 |       3024.06 ± 8.17 |
| phi3 3B Q4_K - Medium          |   2.23 GiB |     3.82 B | SYCL       |  99 |       5 |  none |    0 |           tg128 |         54.19 ± 0.05 |

build: f7e7d2a5 (5331)

### LOG for different GPUs on Linux

In this section there are many logs about this patch working on Linux without affecting performance and or correctness.

##### Lunar Lake
[lnl-test.txt](https://github.com/user-attachments/files/20244936/lnl-test.txt)

[lnl_bench.txt](https://github.com/user-attachments/files/20244938/lnl_bench.txt)
[master_lnl.txt](https://github.com/user-attachments/files/20244973/master_lnl.txt)

##### Battlemage B580
[bmg-test.txt](https://github.com/user-attachments/files/20244863/bmg-test.txt)

[bmg_bench.txt](https://github.com/user-attachments/files/20244885/bmg_bench.txt)
[master_bmg.txt](https://github.com/user-attachments/files/20244865/master_bmg.txt)

##### PVC
[pvc-test.txt](https://github.com/user-attachments/files/20244886/pvc-test.txt)

[pvc_bench.txt](https://github.com/user-attachments/files/20244889/pvc_bench.txt)
[master_pvc.txt](https://github.com/user-attachments/files/20244892/master_pvc.txt)

##### ARC A770
[arc-test.txt](https://github.com/user-attachments/files/20244905/arc-test.txt)

[arc_bench.txt](https://github.com/user-attachments/files/20244915/arc_bench.txt)
[master_arc.txt](https://github.com/user-attachments/files/20244917/master_arc.txt)

#### llama-cli output

[bmg_cli_output.txt](https://github.com/user-attachments/files/20245009/bmg_cli_output.txt)
[lnl_cli_output.txt](https://github.com/user-attachments/files/20245010/lnl_cli_output.txt)
[pvc_cli_output.txt](https://github.com/user-attachments/files/20245012/pvc_cli_output.txt)
[arc_cli_output.txt](https://github.com/user-attachments/files/20245051/arc_cli_output.txt)



Changed files:
filename: ggml/src/ggml-sycl/ggml-sycl.cpp
patch: @@ -385,16 +385,17 @@ static void ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer_t buffer,
     ggml_backend_sycl_buffer_context * ctx = ( ggml_backend_sycl_buffer_context *)buffer->context;
     ggml_sycl_set_device(ctx->device);
     auto stream = &(dpct::dev_mgr::instance().get_device(ctx->device).default_queue());
-    SYCL_CHECK(
-        CHECK_TRY_ERROR(dpct::dev_mgr::instance().get_device(ctx->device).queues_wait_and_throw()));
+    SYCL_CHECK(CHECK_TRY_ERROR(dpct::dev_mgr::instance().get_device(ctx->device).queues_wait_and_throw()));
+#ifndef _WIN32
     // Note: Use host buffer to save the data from mmap(), then copy to device. It's workaround for mmap() issue on PVC GPU.
     // This function will be called during load model from disk. Use memory buffer replace dynamic won't save more time and brings potential memory leak risk here.
-    char* host_buf = (char*)malloc(size);
+    char * host_buf = (char *) malloc(size);
     memcpy(host_buf, data, size);
-    SYCL_CHECK(
-        CHECK_TRY_ERROR((*stream).memcpy((char *)tensor->data + offset, host_buf, size)
-                             .wait()));
+    SYCL_CHECK(CHECK_TRY_ERROR((*stream).memcpy((char *) tensor->data + offset, host_buf, size).wait()));
     free(host_buf);
+#else
+    SYCL_CHECK(CHECK_TRY_ERROR((*stream).memcpy((char *) tensor->data + offset, data, size).wait()));
+#endif
 }
 catch (sycl::exception const &exc) {
   std::cerr << exc.what() << "Exception caught at file:" << __FILE__

filename: tools/llama-bench/README.md
patch: @@ -80,10 +80,6 @@ Using the `-d <n>` option, each test can be run at a specified context depth, pr
 
 For a description of the other options, see the [main example](../main/README.md).
 
-Note:
-
-- When using SYCL backend, there would be hang issue in some cases. Please set `--mmp 0`.
-
 ## Examples
 
 ### Text generation with different models