Compare commits

..

27 Commits

Author SHA1 Message Date
asf0 33ca0dcb9d ggml-hip : add -fno-finite-math-only alongside -ffast-math (#25373)
-ffast-math implies -ffinite-math-only under ROCm/clang 22, which
disables INFINITY/NaN and triggers -Wnan-infinity-disabled (errors
under -Werror in CI). Re-enable infinity handling without dropping
the rest of fast-math.

Fixes #25361
2026-07-07 13:27:50 +02:00
Aman Gupta 024c46ae4e llama: fix quantized kv-cache for dsv4 (#25202) 2026-07-07 17:46:57 +08:00
Neo Zhang 108f186d17 [SYCL] fix unsupported UT cases of CONT & CPY (#25231)
* fix unsupported UT cases of CONT & CPY

* update ops.md

* rm unused head file
2026-07-07 12:20:52 +03:00
Neo Zhang 47e1de77aa [SYCL] support op col2im_1d (#25264)
* support op col2im_1d

* update ops.md

* rm unused words

* update for bf16

* optimize 1%-11% as the review comments

* fix the format issue

* update as the review comments
2026-07-07 11:07:46 +03:00
Neo Zhang 55edb2de44 [SYCL] support OP cross_entropy_loss, cross_entropy_loss_back (#25236)
* support OP cross_entropy_loss, cross_entropy_loss_back

* correct format issue
2026-07-07 10:48:50 +03:00
Todd Malsbary d209086157 sycl : set K_QUANTS_PER_ITERATION to 1 on DMMV path (#25063)
* sycl: add supported types to ggml_sycl_supports_reorder_dmmv

The reordered feature is implemented in ggml_sycl_op_dequantize_mul_mat_vec,
but gated by ggml_sycl_supports_reorder_dmmv. This commit fixes the gate.

Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>

* sycl: set K_QUANTS_PER_ITERATION=1 to improve utilization

When combined with opening the reorder gate, this improves GPU
utilization on B70, giving a significant boost to tg t/s.

Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>

* sycl: replace QK_WARP_SIZE with WARP_SIZE for QK_5

Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>

* sycl: add missing types to ggml_backend_sycl_buffer_init_tensor

Without this, the extra field is not allocated and the reorder path
will not take effect.

Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>

---------

Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
2026-07-07 10:43:41 +03:00
Neo Zhang 95e5254c0a [SYCL] fix unsupport ACC UT cases for noncontiguous (#25124)
* fix unsupport ACC UT cases for noncontiguous

* update ops.md
2026-07-07 10:40:38 +03:00
Neo Zhang 9e5ef0dbb1 sycl : enhance argsort to support all UT cases (#25125) 2026-07-07 10:39:29 +03:00
Neo Zhang 3d4cbdf18a sycl : use sycl func to fix AOT double type issue (#25081) 2026-07-07 10:38:33 +03:00
Neo Zhang 26145b3db7 sycl : rename the env vars from "disable" to "enable" (#25042) 2026-07-07 10:33:51 +03:00
An Long 1a7c25bfdb ggml : make ggml_time_init idempotent (#24422) 2026-07-07 10:29:17 +03:00
o7si defa95c306 speculative : fix out-of-bounds read in ngram-map on prompt shrink (#23936)
* speculative : fix out-of-bounds read in ngram-map on prompt shrink

* speculative : fix ngram-map cleanup cutoff after prompt shrink
2026-07-07 10:25:04 +03:00
fairydreaming a8cfdbb9e4 vulkan : check src0 type in GGML_OP_SET_ROWS to avoid failures due to unimplemented f16 support (#25351)
* vulkan : check src0 type in GGML_OP_SET_ROWS to avoid failures due to unimplemented f16 support

* chore : get rid of else

---------

Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-07-07 12:56:02 +08:00
Hongqiang Wang 6f8895feec opencl: general flash attention decode performance optimizations (#25366)
* opencl: vec flash-attention decode kernels for f16/q8_0/q4_0 KV

* opencl: improve non FA KQ mv kernels

* opencl: tweaks for multiquery FA

* opencl: some tweaks for FA q1 kernels

* opencl: FA with DK=DV=512 for gemma-4

* opencl: various fixes

* opencl: cleanup

* opencl: fix FA decode crash for DK=512 (gemma-4)

The DK=512 decode-only program does not create the f32_f16 prefill
kernel, so the compiled check in ensure_fa_variant never hit and
supports_op gave inconsistent answers for the same op. block_n is also
unset for DK=512 decode; guard it to avoid an out-of-range read at
dispatch.

* opencl: run DK=512 FA decode on CPU

DK=512 decode is bandwidth-bound and faster on the CPU than the GPU,
increasingly so with depth. Decline it in supports_op; prefill stays on the GPU.

* opencl: compile MQ_GQA=8 FA kernels in a minimal program

The full program compiled with -D MQ_GQA=8 runs the Adreno compiler out
of memory at DK>=256. Only the vec_mq kernels are used from this
program, so compile it with FA_MQ_ONLY, which excludes everything else.
Also include the program name in the compile error log.

* opencl: remove stray token in flash_attn_f32_f16.cl

A stray "." broke the f32_f16 program build.

* opencl: split f16-KV FA decode finer (FD_KV_PER_SPLIT_F16)

The 2048 default under-fills the GPU on single-query f16-KV decode;
use 512 for f16 KV to get more splits. Quantized KV keeps 2048.

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-07-06 19:57:52 -07:00
shalinib-ibm ee445f93d8 common: Set optimal default thread count for ppc ( linux as well as AIX) (#25237) 2026-07-07 05:35:20 +08:00
Pascal f36e5c348b metal: add col2im_1d op (f32/f16/bf16) (#25176)
* metal: add col2im_1d op (f32/f16/bf16)

Gather kernel mirroring the CPU/CUDA path: each output (t_out, oc)
reads its ceil(K/s0) source columns with an F32 accumulator, a single
write and no atomics. One thread per output element, 256 per
threadgroup.

* metal: check dst contiguity and type match in supports_op for COL2IM_1D

Align the GGML_OP_COL2IM_1D predicate with the CPU, CUDA, and Vulkan
backends: the kernel writes dst with linear indexing and assumes the
same type as src0, so supports_op must also require a contiguous dst
and op->type == op->src[0]->type.

* Update ggml/src/ggml-metal/ggml-metal.metal

Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>

---------

Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
2026-07-06 20:47:36 +02:00
Johannes Gäßler 74976e1aef CUDA: remove -sm row, refactor cuBLAS (#24216)
* CUDA: remove -sm row, refactor cuBLAS

* fix CDNA + BF16 logic

* fix bad return

* fix src0 strides, contiguous requirements

* fix GGML_CUDA_FORCE_CUBLAS

* fix casts to BF16
2026-07-06 20:04:53 +02:00
Pascal 9abce7473a server: fix deadlock in load_models() when erasing a finished download (#25358)
* server: fix deadlock in load_models() when erasing a finished download

The download monitoring thread acquires the models mutex on its way out,
but load_models() joined it from the erase loop while holding that mutex.
Join it outside the lock via threads_to_join like the other monitoring
threads.

* server: add default timeout to test requests

A hung server now fails the test after 10 minutes instead of stalling
the CI job for hours. Explicit timeouts are unchanged.
2026-07-06 19:26:06 +02:00
Alexey Kopytko cb295bf596 CUDA: extend K-type validation to V-types for flash attention (#24403)
* CUDA: extend K-type validation to V-types for flash attention

* reorder
2026-07-06 16:26:50 +02:00
Xuan-Son Nguyen bfdf581b8b server: temporary skip model downloading API test (#25355) 2026-07-06 16:10:04 +02:00
ragz4125 20a04b2206 ggml-cpu: use UE4M3 LUT in ARM NVFP4 dot product (#25331) 2026-07-06 19:06:40 +08:00
shalinib-ibm 3b4fca11ac ggml-cpu: Enable tiled matmul on AIX (#25199)
The matmul_tiled path uses large local stack buffers for A_pack and B_pack. On AIX this can trigger a segmentation fault, so reduce the buffer footprint there to keep the tiled path usable.

 Performance Impact:
    ~ 2x gains in PP_Speed for FP32, Q4_0 and Q8_0 models tested with llama-bench, llama-batched-bench and llama-cli.
    Models used: Llama3.2 3b Instruct F32, qwen 2.5 3b Q4_0 and Q8_0
2026-07-06 18:18:17 +08:00
hokanosekai 86961efd56 vulkan: fix 32-bit integer overflow in CEIL_DIV (#25245) 2026-07-06 10:35:57 +02:00
Pascal d80e878501 ui: restore Ctrl+B sidebar toggle shortcut (#25307) 2026-07-06 10:30:07 +02:00
Adrien Gallouët 48719618e8 scripts : use HF_TOKEN when downloading UI assets (#25280)
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-07-06 09:53:35 +02:00
a-huk d06ddd3589 ggml-hip: enable -ffast-math for HIP builds (#23862) 2026-07-06 15:02:26 +08:00
Xuan-Son Nguyen 898b08854d ui: fake 200 for proxy DELETE req (#25298) 2026-07-06 08:41:39 +02:00
57 changed files with 7221 additions and 2111 deletions
+22 -1
View File
@@ -55,6 +55,10 @@
#include <pwd.h>
#endif
#if defined(_AIX)
#include <sys/systemcfg.h>
#endif
#if defined(_MSC_VER)
#pragma warning(disable: 4244 4267) // possible loss of data
#endif
@@ -72,7 +76,16 @@ common_time_meas::~common_time_meas() {
//
int32_t common_cpu_get_num_physical_cores() {
#ifdef __linux__
#if defined(_AIX)
int32_t logical_cpus = _system_configuration.ncpus;
int32_t smt_threads = _system_configuration.smt_threads;
if (smt_threads > 0) {
return static_cast<int32_t>(logical_cpus / smt_threads);
}
if (logical_cpus > 0) {
return static_cast<int32_t>(logical_cpus);
}
#elif defined(__linux__)
// enumerate the set of thread siblings, num entries is num cores
std::unordered_set<std::string> siblings;
for (uint32_t cpu=0; cpu < UINT32_MAX; ++cpu) {
@@ -202,6 +215,14 @@ int32_t common_cpu_get_num_math() {
}
}
}
#elif defined(__powerpc64__) || defined(__powerpc__)
int32_t smt_factor = 1;
int phy_cpus = common_cpu_get_num_physical_cores();
int logical_cpus = sysconf(_SC_NPROCESSORS_ONLN);
if (phy_cpus > 0 && logical_cpus > phy_cpus) {
smt_factor = logical_cpus / phy_cpus;
}
return phy_cpus * std::min(smt_factor, 2);
#endif
return common_cpu_get_num_physical_cores();
}
+15 -9
View File
@@ -125,6 +125,16 @@ void common_ngram_map_begin(
LOG_DBG("%s: begin, idx_last_draft=%zu, new begin=%zu, #keys=%zu\n", __func__,
map.idx_last_check, size_begin, map.keys.size());
size_t idx_begin_cleanup = map.size_last_begin;
if (idx_begin_cleanup > size_begin) {
if (size_begin > (size_t) map.size_key + map.size_value) {
idx_begin_cleanup = size_begin - map.size_key - map.size_value;
} else {
idx_begin_cleanup = 0;
}
LOG_INF("%s: shrink cleanup begin: %zu -> %zu\n", __func__, map.size_last_begin, idx_begin_cleanup);
}
size_t count_map_entries_upd = 0;
if (!map.key_map.empty() && size_begin < map.idx_last_check) {
if (map.show_key_map_stats) {
@@ -150,27 +160,23 @@ void common_ngram_map_begin(
// Update the map from hash to key index (clear outdated entries).
for (size_t i = 0; i < map.key_map.size(); ++i) {
uint32_t key_idx = map.key_map[i];
if (key_idx >= map.size_last_begin) {
if (key_idx != 0 && key_idx >= idx_begin_cleanup) {
map.key_map[i] = 0;
count_map_entries_upd++;
}
}
map.key_map_last_idx = (map.size_last_begin > 0) ? map.size_last_begin - 1 : 0;
map.key_map_last_idx = (idx_begin_cleanup > 0) ? (uint32_t) (idx_begin_cleanup - 1) : 0;
}
if (size_begin < map.idx_last_check && !map.keys.empty()) {
// The next token generation will start at index size_begin.
// The tokens between map.size_last_begin and size_begin are no longer valid.
//
// Refresh map: Remove all entries with index >= map.size_last_begin.
size_t count_keys = map.keys.size();
size_t count_keys_del = 0;
size_t count_values_del = 0;
for (int32_t i = map.keys.size() - 1; i >= 0; --i) {
common_ngram_map_key & key = map.keys[i];
if (key.key_idx >= map.size_last_begin) {
if (key.key_idx >= idx_begin_cleanup) {
// Delete the key.
LOG_DBG("%s: delete key %d at index %zu (>= size_last_begin=%zu)\n", __func__, i, key.key_idx, map.size_last_begin);
LOG_DBG("%s: delete key %d at index %zu (>= idx_begin_cleanup=%zu)\n", __func__, i, key.key_idx, idx_begin_cleanup);
map.keys.erase(map.keys.begin() + i);
count_keys_del++;
continue;
@@ -182,7 +188,7 @@ void common_ngram_map_begin(
// Check the indices of the values.
for (int16_t j = COMMON_NGRAM_MAX_VALUES - 1; j >= 0; --j) {
common_ngram_map_value & value = key.values[j];
if (value.value_idx >= map.size_last_begin) {
if (value.value_idx != 0 && value.value_idx >= idx_begin_cleanup) {
// Delete the value.
count_values_del++;
+4 -4
View File
@@ -790,10 +790,10 @@ use 1 SYCL GPUs: [0] with Max compute units:512
| GGML_SYCL_DEBUG | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG |
| GGML_SYCL_DEV2DEV_MEMCPY | 0 (default) or 1 | Choose the SYCL or L0 API in dev2dev memory copy.<br>Value: <br>* 0: SYCL API (default)<br>* 1: L0 API -- L0 API is found to lead to abnormal crash in some case. This debug flag is used to check the issue.|
| GGML_SYCL_ENABLE_FLASH_ATTN | 1 (default) or 0| Enable Flash-Attention. It can reduce memory usage. The performance impact depends on the LLM.|
| 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 SYCL Graph is still on development, no better performance. |
| GGML_SYCL_ENABLE_OPT | 0 or 1 (default)| Enable optimize features for Intel GPUs. (Recommended to 0 for Intel devices older than Gen 10) |
| GGML_SYCL_ENABLE_GRAPH | 0 (default) or 1 | Enable running computations through SYCL Graphs feature. Disabled by default because SYCL Graph is still on development, no better performance. |
| GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).|
| GGML_SYCL_DISABLE_DNN | 0 (default) or 1 | Disable running computations through oneDNN and always use oneMKL. |
| GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. |
| GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. |
| 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 |
| UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS | 0 (default) or 1 | Allow SYCL/Unified Runtime Level Zero device allocations larger than 4 GiB. llama.cpp's direct Level Zero allocation path requests the relaxed maximum-size limit itself when GGML_SYCL_ENABLE_LEVEL_ZERO=1. |
@@ -807,7 +807,7 @@ Pass these via `CXXFLAGS` or add a one-off `#define` to enable a flag on the spo
|-----------------|----------------------------------------------------------------------------------|
| DEBUG_SYCL_POOL | Enable device memory pool logging on teardown. Useful for profiling allocations. |
| DEBUG_SYCL_MALLOC | Enable verbose per-call logging of device pool alloc/free operations. |
| GGML_SYCL_SUPPORT_VMM | Support to building with VMM code. Default is Yes. |
## Design Rule
+3 -6
View File
@@ -270,13 +270,10 @@ The environment variable [`CUDA_SCALE_LAUNCH_QUEUES`](https://docs.nvidia.com/cu
Consider setting `CUDA_SCALE_LAUNCH_QUEUES=4x`, which increases the CUDA command buffer to 4 times its default size. This optimization is particularly beneficial for **Multi-GPU setups with pipeline parallelism**, where it significantly improves prompt processing throughput by allowing more operations to be enqueued across GPUs.
#### GGML_CUDA_FORCE_CUBLAS_COMPUTE_32F
#### GGML_CUDA_CUBLAS_COMPUTE_TYPE
Use `GGML_CUDA_FORCE_CUBLAS_COMPUTE_32F` environment variable to use FP32 compute type on all GPUs in FP16 cuBLAS for preventing possible numerical overflows in exchange for slower prompt processing (small impact on RTX PRO/Datacenter products and significant on GeForce products).
#### GGML_CUDA_FORCE_CUBLAS_COMPUTE_16F
Use `GGML_CUDA_FORCE_CUBLAS_COMPUTE_16F` environment variable to force use FP16 compute type (instead of default FP32) in FP16 cuBLAS for V100, CDNA and RDNA4.
Override default, speed-optimized compute types for cuBLAS matrix multiplications.
Legal values: `auto`, `f16`, `fp16`, `bf16`, `f32`, `fp32`.
### Unified Memory
+6 -6
View File
@@ -21,12 +21,12 @@ Legend:
| ADD_ID | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
| ARANGE | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| ARGMAX | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| ARGSORT | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ✅ | ❌ | ❌ |
| ARGSORT | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | | ✅ | ✅ | ❌ | ❌ |
| CEIL | ❌ | ❌ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| CLAMP | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ❌ | ❌ |
| COL2IM_1D | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| COL2IM_1D | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| CONCAT | ❌ | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| CONT | ❌ | 🟡 | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ❌ | ❌ |
| CONT | ❌ | 🟡 | ✅ | ✅ | ✅ | 🟡 | | ✅ | 🟡 | ❌ | ❌ |
| CONV_2D | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
| CONV_2D_DW | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| CONV_3D | ❌ | ❌ | ✅ | ❌ | ✅ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
@@ -35,8 +35,8 @@ Legend:
| COS | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | 🟡 | ✅ | ❌ | ❌ |
| COUNT_EQUAL | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| CPY | ❌ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ❌ | ❌ |
| CROSS_ENTROPY_LOSS | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| CROSS_ENTROPY_LOSS_BACK | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| CROSS_ENTROPY_LOSS | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| CROSS_ENTROPY_LOSS_BACK | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | | ❌ | ❌ | ❌ | ❌ |
| CUMSUM | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| DIAG | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| DIAG_MASK_INF | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
@@ -70,7 +70,7 @@ Legend:
| MUL | ❌ | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
| MUL_MAT | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 |
| MUL_MAT_HADAMARD | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| MUL_MAT_ID | ❌ | 🟡 | ✅ | ✅ | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | 🟡 | ❌ |
| MUL_MAT_ID | ❌ | 🟡 | ✅ | ✅ | 🟡 | 🟡 | | ✅ | 🟡 | 🟡 | ❌ |
| NEG | ❌ | ✅ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| NORM | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | ❌ | ❌ |
| OPT_STEP_ADAMW | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ |
+555 -471
View File
File diff suppressed because it is too large Load Diff
-3
View File
@@ -30,9 +30,6 @@ GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_buffer_type(int de
// conduct allreduce operation between devices
GGML_BACKEND_API bool ggml_backend_cuda_allreduce_tensor(ggml_backend_t * backends, struct ggml_tensor ** tensors, size_t n_backends);
// split tensor buffer that splits matrices by rows across multiple devices
GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_split_buffer_type(int main_device, const float * tensor_split);
// pinned host buffer for use with the CPU backend for faster copies between CPU and GPU
GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_host_buffer_type(void);
+4 -4
View File
@@ -812,10 +812,10 @@ void ggml_vec_dot_nvfp4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const vo
const float dy0 = GGML_CPU_FP16_TO_FP32(y[2*ib].d);
const float dy1 = GGML_CPU_FP16_TO_FP32(y[2*ib+1].d);
const float32x4_t nvsc = {
ggml_ue4m3_to_fp32(x[ib].d[0]),
ggml_ue4m3_to_fp32(x[ib].d[1]),
ggml_ue4m3_to_fp32(x[ib].d[2]),
ggml_ue4m3_to_fp32(x[ib].d[3])
GGML_CPU_UE4M3_TO_FP32(x[ib].d[0]),
GGML_CPU_UE4M3_TO_FP32(x[ib].d[1]),
GGML_CPU_UE4M3_TO_FP32(x[ib].d[2]),
GGML_CPU_UE4M3_TO_FP32(x[ib].d[3])
};
const float32x4_t scales = vmulq_f32(nvsc, (float32x4_t){dy0, dy0, dy1, dy1});
+20 -14
View File
@@ -2321,24 +2321,28 @@ class tinyBLAS_Q0_PPC {
}
void matmul(int64_t m, int64_t n) {
#if defined(_AIX) || defined(__BIG_ENDIAN__)
mnpack(0, m, 0, n);
#else
const int64_t mc = 64;
const int64_t kc = 64;
int64_t mc = 64;
int64_t nc = 64;
int64_t kc = 64;
int64_t n_chunk = 64;
#if defined(_AIX) || defined(__BIG_ENDIAN__)
mc = 32;
nc = 32;
kc = 32;
n_chunk = 32
#endif
int64_t n_aligned = 0;
if (n % 64 == 0) {
if (n % n_chunk == 0) {
n_aligned = n;
} else if (n == 4) {
n_aligned = 4;
} else if (n < 64) {
} else if (n < n_chunk) {
n_aligned = (n / 8) * 8;
} else {
n_aligned = (n / 64) * 64;
n_aligned = (n / n_chunk) * n_chunk;
}
if (n_aligned > 0) {
if (n_aligned % 64 == 0) nc = 64;
if (n_aligned % n_chunk == 0) nc = n_chunk;
else if (n_aligned == n) nc = n;
else if (n_aligned % 32 == 0) nc = 32;
else if (n_aligned % 24 == 0) nc = 24;
@@ -2354,7 +2358,6 @@ class tinyBLAS_Q0_PPC {
} else {
mnpack(0, m, 0, n);
}
#endif
}
private:
@@ -3195,16 +3198,19 @@ class tinyBLAS_PPC {
}
void matmul(int64_t m, int64_t n) {
int64_t mc = 256;
int64_t nc = 256;
int64_t kc = 256;
#if defined(_AIX) || defined(__BIG_ENDIAN__)
mnpack(0, m, 0, n);
#else
int64_t mc = 256; int64_t nc = 256; int64_t kc = 256;
mc = 128;
nc = 128;
kc = 128;
#endif
if (m % mc == 0 && n % nc == 0 && k % kc == 0) {
matmul_tiled(m, n, mc, nc, kc);
} else {
mnpack(0, m, 0, n);
}
#endif
}
private:
+2 -2
View File
@@ -131,8 +131,8 @@ extern float ggml_table_f32_ue4m3[1 << 8];
#define GGML_CPU_E8M0_TO_FP32_HALF(x) GGML_E8M0_TO_FP32_HALF(x)
#endif
// Use lookup table for UE4M3 on x86 (faster than bit manipulation)
#if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__)
// Use lookup table for UE4M3 on x86 and ARM (faster than bit manipulation)
#if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) || defined(__ARM_NEON)
#define GGML_CPU_UE4M3_TO_FP32(x) ggml_table_f32_ue4m3[(uint8_t)(x)]
#else
#define GGML_CPU_UE4M3_TO_FP32(x) ggml_ue4m3_to_fp32(x)
+86 -34
View File
@@ -104,8 +104,8 @@ static __global__ void dequantize_block_q4_0(const void * __restrict__ vx, dst_t
const uint8_t * q = x->qs + 4*il;
for (int l = 0; l < 4; ++l) {
y[l+ 0] = d * (q[l] & 0xF) + dm;
y[l+16] = d * (q[l] >> 4) + dm;
y[l+ 0] = ggml_cuda_cast<dst_t>(d * (q[l] & 0xF) + dm);
y[l+16] = ggml_cuda_cast<dst_t>(d * (q[l] >> 4) + dm);
}
}
@@ -131,8 +131,8 @@ static __global__ void dequantize_block_q4_1(const void * __restrict__ vx, dst_t
const uint8_t * q = x->qs + 4*il;
for (int l = 0; l < 4; ++l) {
y[l+ 0] = d.x * (q[l] & 0xF) + d.y;
y[l+16] = d.x * (q[l] >> 4) + d.y;
y[l+ 0] = ggml_cuda_cast<dst_t>(d.x * (q[l] & 0xF) + d.y);
y[l+16] = ggml_cuda_cast<dst_t>(d.x * (q[l] >> 4) + d.y);
}
}
@@ -154,10 +154,10 @@ static __global__ void dequantize_block_q2_K(const void * __restrict__ vx, dst_t
float dall = __low2half(x[i].dm);
float dmin = __high2half(x[i].dm);
y[l+ 0] = dall * (x[i].scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (x[i].scales[is+0] >> 4);
y[l+32] = dall * (x[i].scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (x[i].scales[is+2] >> 4);
y[l+64] = dall * (x[i].scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (x[i].scales[is+4] >> 4);
y[l+96] = dall * (x[i].scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (x[i].scales[is+6] >> 4);
y[l+ 0] = ggml_cuda_cast<dst_t>(dall * (x[i].scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (x[i].scales[is+0] >> 4));
y[l+32] = ggml_cuda_cast<dst_t>(dall * (x[i].scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (x[i].scales[is+2] >> 4));
y[l+64] = ggml_cuda_cast<dst_t>(dall * (x[i].scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (x[i].scales[is+4] >> 4));
y[l+96] = ggml_cuda_cast<dst_t>(dall * (x[i].scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (x[i].scales[is+6] >> 4));
}
template<typename dst_t>
@@ -188,7 +188,9 @@ static __global__ void dequantize_block_q3_K(const void * __restrict__ vx, dst_t
const uint8_t * q = x[i].qs + 32*n;
const uint8_t * hm = x[i].hmask;
for (int l = l0; l < l0+4; ++l) y[l] = dl * ((int8_t)((q[l] >> shift) & 3) - ((hm[l] & m) ? 0 : 4));
for (int l = l0; l < l0+4; ++l) {
y[l] = ggml_cuda_cast<dst_t>(dl * ((int8_t)((q[l] >> shift) & 3) - ((hm[l] & m) ? 0 : 4)));
}
}
static inline __device__ void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m) {
@@ -226,8 +228,8 @@ static __global__ void dequantize_block_q4_K(const void * __restrict__ vx, dst_t
get_scale_min_k4(is + 1, x[i].scales, sc, m);
const float d2 = dall * sc; const float m2 = dmin * m;
for (int l = 0; l < n; ++l) {
y[l + 0] = d1 * (q[l] & 0xF) - m1;
y[l +32] = d2 * (q[l] >> 4) - m2;
y[l + 0] = ggml_cuda_cast<dst_t>(d1 * (q[l] & 0xF) - m1);
y[l +32] = ggml_cuda_cast<dst_t>(d2 * (q[l] >> 4) - m2);
}
}
@@ -258,11 +260,11 @@ static __global__ void dequantize_block_q5_K(const void * __restrict__ vx, dst_t
const float d2 = dall * sc; const float m2 = dmin * m;
uint8_t hm = 1 << (2*il);
y[ 0] = d1 * ((ql[ 0] & 0xF) + (qh[ 0] & hm ? 16 : 0)) - m1;
y[ 1] = d1 * ((ql[ 1] & 0xF) + (qh[ 1] & hm ? 16 : 0)) - m1;
y[ 0] = ggml_cuda_cast<dst_t>(d1 * ((ql[ 0] & 0xF) + (qh[ 0] & hm ? 16 : 0)) - m1);
y[ 1] = ggml_cuda_cast<dst_t>(d1 * ((ql[ 1] & 0xF) + (qh[ 1] & hm ? 16 : 0)) - m1);
hm <<= 1;
y[32] = d2 * ((ql[ 0] >> 4) + (qh[ 0] & hm ? 16 : 0)) - m2;
y[33] = d2 * ((ql[ 1] >> 4) + (qh[ 1] & hm ? 16 : 0)) - m2;
y[32] = ggml_cuda_cast<dst_t>(d2 * ((ql[ 0] >> 4) + (qh[ 0] & hm ? 16 : 0)) - m2);
y[33] = ggml_cuda_cast<dst_t>(d2 * ((ql[ 1] >> 4) + (qh[ 1] & hm ? 16 : 0)) - m2);
}
template<typename dst_t>
@@ -285,10 +287,10 @@ static __global__ void dequantize_block_q6_K(const void * __restrict__ vx, dst_t
const uint8_t qh = x[i].qh[32*ip + il];
const int8_t * sc = x[i].scales + is;
y[ 0] = d * sc[0] * ((int8_t)((ql[ 0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32);
y[32] = d * sc[2] * ((int8_t)((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32);
y[64] = d * sc[4] * ((int8_t)((ql[ 0] >> 4) | (((qh >> 4) & 3) << 4)) - 32);
y[96] = d * sc[6] * ((int8_t)((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32);
y[ 0] = ggml_cuda_cast<dst_t>(d * sc[0] * ((int8_t)((ql[ 0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32));
y[32] = ggml_cuda_cast<dst_t>(d * sc[2] * ((int8_t)((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32));
y[64] = ggml_cuda_cast<dst_t>(d * sc[4] * ((int8_t)((ql[ 0] >> 4) | (((qh >> 4) & 3) << 4)) - 32));
y[96] = ggml_cuda_cast<dst_t>(d * sc[6] * ((int8_t)((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32));
}
template<typename dst_t>
@@ -307,7 +309,9 @@ static __global__ void dequantize_block_iq2_xxs(const void * __restrict__ vx, ds
const uint32_t aux32 = q2[2] | (q2[3] << 16);
const float d = (float)x[i].d * (0.5f + (aux32 >> 28)) * 0.25f;
const uint8_t signs = ksigns_iq2xs[(aux32 >> 7*il) & 127];
for (int j = 0; j < 8; ++j) y[j] = d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f);
for (int j = 0; j < 8; ++j) {
y[j] = ggml_cuda_cast<dst_t>(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f));
}
}
template<typename dst_t>
@@ -324,7 +328,9 @@ static __global__ void dequantize_block_iq2_xs(const void * __restrict__ vx, dst
const uint8_t * grid = (const uint8_t *)(iq2xs_grid + (q2[il] & 511));
const float d = (float)x[i].d * (0.5f + ((x[i].scales[ib] >> 4*(il/2)) & 0xf)) * 0.25f;
const uint8_t signs = ksigns_iq2xs[q2[il] >> 9];
for (int j = 0; j < 8; ++j) y[j] = d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f);
for (int j = 0; j < 8; ++j) {
y[j] = ggml_cuda_cast<dst_t>(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f));
}
}
template<typename dst_t>
@@ -340,7 +346,9 @@ static __global__ void dequantize_block_iq2_s(const void * __restrict__ vx, dst_
const uint8_t * grid = (const uint8_t *)(iq2s_grid + (x[i].qs[4*ib+il] | ((x[i].qh[ib] << (8-2*il)) & 0x300)));
const float d = (float)x[i].d * (0.5f + ((x[i].scales[ib] >> 4*(il/2)) & 0xf)) * 0.25f;
const uint8_t signs = x[i].qs[QK_K/8+4*ib+il];
for (int j = 0; j < 8; ++j) y[j] = d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f);
for (int j = 0; j < 8; ++j) {
y[j] = ggml_cuda_cast<dst_t>(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f));
}
}
template<typename dst_t>
@@ -361,8 +369,8 @@ static __global__ void dequantize_block_iq3_xxs(const void * __restrict__ vx, ds
const float d = (float)x[i].d * (0.5f + (aux32 >> 28)) * 0.5f;
const uint8_t signs = ksigns_iq2xs[(aux32 >> 7*il) & 127];
for (int j = 0; j < 4; ++j) {
y[j+0] = d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f);
y[j+4] = d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f);
y[j+0] = ggml_cuda_cast<dst_t>(d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f));
y[j+4] = ggml_cuda_cast<dst_t>(d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f));
}
}
@@ -382,8 +390,8 @@ static __global__ void dequantize_block_iq3_s(const void * __restrict__ vx, dst_
const float d = (float)x[i].d * (1 + 2*((x[i].scales[ib/2] >> 4*(ib%2)) & 0xf));
const uint8_t signs = x[i].signs[4*ib + il];
for (int j = 0; j < 4; ++j) {
y[j+0] = d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f);
y[j+4] = d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f);
y[j+0] = ggml_cuda_cast<dst_t>(d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f));
y[j+4] = ggml_cuda_cast<dst_t>(d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f));
}
}
@@ -404,7 +412,7 @@ static __global__ void dequantize_block_iq1_s(const void * __restrict__ vx, dst_
grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f;
grid32[0] &= 0x0f0f0f0f;
for (int j = 0; j < 8; ++j) {
y[j] = d * (q[j] + delta);
y[j] = ggml_cuda_cast<dst_t>(d * (q[j] + delta));
}
}
@@ -429,7 +437,7 @@ static __global__ void dequantize_block_iq1_m(const void * __restrict__ vx, dst_
grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f;
grid32[0] &= 0x0f0f0f0f;
for (int j = 0; j < 8; ++j) {
y[j] = d * (q[j] + delta);
y[j] = ggml_cuda_cast<dst_t>(d * (q[j] + delta));
}
}
@@ -446,8 +454,8 @@ static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst
const uint8_t * q4 = x[ib].qs + 4*il;
const float d = (float)x[ib].d;
for (int j = 0; j < 4; ++j) {
y[j+ 0] = d * kvalues_iq4nl[q4[j] & 0xf];
y[j+16] = d * kvalues_iq4nl[q4[j] >> 4];
y[j+ 0] = ggml_cuda_cast<dst_t>(d * kvalues_iq4nl[q4[j] & 0xf]);
y[j+16] = ggml_cuda_cast<dst_t>(d * kvalues_iq4nl[q4[j] >> 4]);
}
}
@@ -463,8 +471,8 @@ static __global__ void dequantize_block_iq4_xs(const void * __restrict__ vx, dst
const uint8_t * q4 = x[i].qs + 16*ib + 4*il;
const float d = (float)x[i].d * ((((x[i].scales_l[ib/2] >> 4*(ib%2)) & 0xf) | (((x[i].scales_h >> 2*ib) & 3) << 4)) - 32);
for (int j = 0; j < 4; ++j) {
y[j+ 0] = d * kvalues_iq4nl[q4[j] & 0xf];
y[j+16] = d * kvalues_iq4nl[q4[j] >> 4];
y[j+ 0] = ggml_cuda_cast<dst_t>(d * kvalues_iq4nl[q4[j] & 0xf]);
y[j+16] = ggml_cuda_cast<dst_t>(d * kvalues_iq4nl[q4[j] >> 4]);
}
}
@@ -481,8 +489,8 @@ static __global__ void dequantize_block_mxfp4(const void * __restrict__ vx, dst_
const uint8_t * q4 = x[ib].qs + 4*il;
const float d = ggml_cuda_e8m0_to_fp32(x[ib].e);
for (int j = 0; j < 4; ++j) {
y[j+ 0] = d * kvalues_mxfp4[q4[j] & 0xf]*0.5f;
y[j+16] = d * kvalues_mxfp4[q4[j] >> 4]*0.5f;
y[j+ 0] = ggml_cuda_cast<dst_t>(d * kvalues_mxfp4[q4[j] & 0xf]*0.5f);
y[j+16] = ggml_cuda_cast<dst_t>(d * kvalues_mxfp4[q4[j] >> 4]*0.5f);
}
}
@@ -700,6 +708,50 @@ static void convert_unary_cont_cuda(const void * vx, dst_t * y, const int64_t k,
to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
switch (type) {
case GGML_TYPE_Q1_0:
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q4_0:
return dequantize_row_q4_0_cuda;
case GGML_TYPE_Q4_1:
return dequantize_row_q4_1_cuda;
case GGML_TYPE_Q5_0:
return dequantize_block_cont_cuda<QK5_0, QR5_0, dequantize_q5_0>;
case GGML_TYPE_Q5_1:
return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cont_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_Q2_K:
return dequantize_row_q2_K_cuda;
case GGML_TYPE_Q3_K:
return dequantize_row_q3_K_cuda;
case GGML_TYPE_Q4_K:
return dequantize_row_q4_K_cuda;
case GGML_TYPE_Q5_K:
return dequantize_row_q5_K_cuda;
case GGML_TYPE_Q6_K:
return dequantize_row_q6_K_cuda;
case GGML_TYPE_IQ2_XXS:
return dequantize_row_iq2_xxs_cuda;
case GGML_TYPE_IQ2_XS:
return dequantize_row_iq2_xs_cuda;
case GGML_TYPE_IQ2_S:
return dequantize_row_iq2_s_cuda;
case GGML_TYPE_IQ3_XXS:
return dequantize_row_iq3_xxs_cuda;
case GGML_TYPE_IQ1_S:
return dequantize_row_iq1_s_cuda;
case GGML_TYPE_IQ1_M:
return dequantize_row_iq1_m_cuda;
case GGML_TYPE_IQ4_NL:
return dequantize_row_iq4_nl_cuda;
case GGML_TYPE_IQ4_XS:
return dequantize_row_iq4_xs_cuda;
case GGML_TYPE_IQ3_S:
return dequantize_row_iq3_s_cuda;
case GGML_TYPE_MXFP4:
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F32:
return convert_unary_cont_cuda<float>;
case GGML_TYPE_F16:
+22 -16
View File
@@ -337,6 +337,26 @@ enum best_fattn_kernel {
BEST_FATTN_KERNEL_MMA_F16 = 400,
};
static bool ggml_cuda_fattn_kv_type_supported(ggml_type type) {
switch (type) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
return true;
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
#ifndef GGML_CUDA_FA_ALL_QUANTS
return false;
#endif // GGML_CUDA_FA_ALL_QUANTS
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q8_0:
case GGML_TYPE_BF16:
return true;
default:
return false;
}
}
static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const ggml_tensor * dst) {
#ifndef FLASH_ATTN_AVAILABLE
GGML_UNUSED(device); GGML_UNUSED(dst);
@@ -427,22 +447,8 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
}
#endif // GGML_CUDA_FA_ALL_QUANTS
switch (K->type) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
break;
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
#ifndef GGML_CUDA_FA_ALL_QUANTS
return BEST_FATTN_KERNEL_NONE;
#endif // GGML_CUDA_FA_ALL_QUANTS
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q8_0:
case GGML_TYPE_BF16:
break;
default:
return BEST_FATTN_KERNEL_NONE;
if (!ggml_cuda_fattn_kv_type_supported(K->type) || !ggml_cuda_fattn_kv_type_supported(V->type)) {
return BEST_FATTN_KERNEL_NONE;
}
if (mask && mask->ne[2] != 1) {
File diff suppressed because it is too large Load Diff
+3
View File
@@ -278,6 +278,9 @@ int get_mmvq_mmid_max_batch(ggml_type type, int cc) {
}
bool ggml_cuda_should_use_mmvq(enum ggml_type type, int cc, int64_t ne11) {
if (!ggml_is_quantized(type)) {
return false;
}
if (GGML_CUDA_CC_IS_CDNA(cc)) {
if (GGML_CUDA_CC_IS_CDNA1(cc)) {
switch (type) {
+2
View File
@@ -155,3 +155,5 @@ if (GGML_HIP_RCCL)
endif()
target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas)
target_compile_options(ggml-hip PRIVATE "$<$<COMPILE_LANGUAGE:HIP>:-ffast-math;-fno-finite-math-only>")
+20
View File
@@ -1800,6 +1800,26 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_1
return res;
}
ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d(ggml_metal_library_t lib, const ggml_tensor * op) {
assert(op->op == GGML_OP_COL2IM_1D);
GGML_ASSERT(ggml_is_contiguous(op->src[0]));
GGML_ASSERT(op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_BF16);
char base[256];
char name[256];
snprintf(base, 256, "kernel_col2im_1d_%s", ggml_type_name(op->src[0]->type));
snprintf(name, 256, "%s", base);
ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
if (!res.pipeline) {
res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr);
}
return res;
}
ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d(ggml_metal_library_t lib, const ggml_tensor * op) {
assert(op->op == GGML_OP_CONV_TRANSPOSE_2D);
+1
View File
@@ -150,6 +150,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_rope
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_im2col (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_1d (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_upscale (ggml_metal_library_t lib, const struct ggml_tensor * op);
+5
View File
@@ -1157,6 +1157,11 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32) &&
op->src[1]->type == GGML_TYPE_F32 &&
op->type == GGML_TYPE_F32;
case GGML_OP_COL2IM_1D:
return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_BF16) &&
op->type == op->src[0]->type &&
ggml_is_contiguous(op->src[0]) &&
ggml_is_contiguous(op);
case GGML_OP_CONV_3D:
return ggml_is_contiguous(op->src[0]) &&
ggml_is_contiguous(op->src[1]) &&
+10
View File
@@ -603,6 +603,16 @@ typedef struct {
uint64_t nb1;
} ggml_metal_kargs_conv_transpose_1d;
typedef struct {
int32_t T_in;
int32_t T_out;
int32_t OC;
int32_t K;
int32_t K_OC;
int32_t s0;
int32_t p0;
} ggml_metal_kargs_col2im_1d;
typedef struct {
int32_t IC;
int32_t IH;
+45
View File
@@ -395,6 +395,10 @@ static int ggml_metal_op_encode_impl(ggml_metal_op_t ctx, int idx) {
{
n_fuse = ggml_metal_op_conv_transpose_2d(ctx, idx);
} break;
case GGML_OP_COL2IM_1D:
{
n_fuse = ggml_metal_op_col2im_1d(ctx, idx);
} break;
case GGML_OP_CONV_3D:
{
n_fuse = ggml_metal_op_conv_3d(ctx, idx);
@@ -3854,6 +3858,47 @@ int ggml_metal_op_conv_transpose_1d(ggml_metal_op_t ctx, int idx) {
return 1;
}
int ggml_metal_op_col2im_1d(ggml_metal_op_t ctx, int idx) {
ggml_tensor * op = ctx->node(idx);
ggml_metal_library_t lib = ctx->lib;
ggml_metal_encoder_t enc = ctx->enc;
const int32_t s0 = ((const int32_t *)(op->op_params))[0];
const int32_t OC = ((const int32_t *)(op->op_params))[1];
const int32_t p0 = ((const int32_t *)(op->op_params))[2];
const int32_t K_OC = (int32_t) op->src[0]->ne[0];
const int32_t T_in = (int32_t) op->src[0]->ne[1];
const int32_t K = K_OC / OC;
const int32_t T_out = (int32_t) op->ne[0];
ggml_metal_kargs_col2im_1d args = {
/*.T_in =*/ T_in,
/*.T_out =*/ T_out,
/*.OC =*/ OC,
/*.K =*/ K,
/*.K_OC =*/ K_OC,
/*.s0 =*/ s0,
/*.p0 =*/ p0,
};
auto pipeline = ggml_metal_library_get_pipeline_col2im_1d(lib, op);
const int total = T_out * OC;
const int nth = 256;
const int ntg = (total + nth - 1) / nth;
ggml_metal_encoder_set_pipeline(enc, pipeline);
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0);
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1);
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 2);
ggml_metal_encoder_dispatch_threadgroups(enc, ntg, 1, 1, nth, 1, 1);
return 1;
}
int ggml_metal_op_conv_transpose_2d(ggml_metal_op_t ctx, int idx) {
ggml_tensor * op = ctx->node(idx);
+1
View File
@@ -78,6 +78,7 @@ int ggml_metal_op_conv_2d (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_conv_3d (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_conv_transpose_1d (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_conv_transpose_2d (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_col2im_1d (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_upscale (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_pad (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_pad_reflect_1d (ggml_metal_op_t ctx, int idx);
+43
View File
@@ -4977,6 +4977,49 @@ kernel void kernel_conv_transpose_1d<half>(
uint3 tgpg[[threadgroups_per_grid]]);
template <typename T>
kernel void kernel_col2im_1d(
constant ggml_metal_kargs_col2im_1d & args,
device const T * col,
device T * dst,
uint tgpig [[threadgroup_position_in_grid]],
uint tpitg [[thread_position_in_threadgroup]],
uint ntg [[threads_per_threadgroup]]) {
const int idx = tgpig * ntg + tpitg;
if (idx >= args.T_out * args.OC) {
return;
}
const int t_out = idx % args.T_out;
const int oc = idx / args.T_out;
const int t_abs = t_out + args.p0; // absolute position in uncropped signal
int t_in_min = (t_abs - args.K + args.s0) / args.s0; // ceil((t_abs - K + 1) / s0)
if (t_in_min < 0) {
t_in_min = 0;
}
int t_in_max = t_abs / args.s0;
if (t_in_max >= args.T_in) {
t_in_max = args.T_in - 1;
}
float sum = 0.0f;
for (int t_in = t_in_min; t_in <= t_in_max; t_in++) {
const int k = t_abs - t_in * args.s0;
sum += float(col[(oc * args.K + k) + t_in * args.K_OC]);
}
dst[t_out + oc * args.T_out] = T(sum);
}
template [[host_name("kernel_col2im_1d_f32")]] kernel void kernel_col2im_1d<float>(constant ggml_metal_kargs_col2im_1d &, device const float *, device float *, uint, uint, uint);
template [[host_name("kernel_col2im_1d_f16")]] kernel void kernel_col2im_1d<half>(constant ggml_metal_kargs_col2im_1d &, device const half *, device half *, uint, uint, uint);
#if defined(GGML_METAL_HAS_BF16)
template [[host_name("kernel_col2im_1d_bf16")]] kernel void kernel_col2im_1d<bfloat>(constant ggml_metal_kargs_col2im_1d &, device const bfloat *, device bfloat *, uint, uint, uint);
#endif
typedef void (conv_transpose_2d_t)(
constant ggml_metal_kargs_conv_transpose_2d & args,
device const float * src0,
+1
View File
@@ -20,6 +20,7 @@ static const ggml_opencl_fa_dim g_fa_dims_adreno_default[] = {
{192, 128, 16, 16, 1, 0},
{192, 192, 16, 16, 1, 0},
{256, 256, 16, 16, 16, 0},
{512, 512, 8, 16, 64, 0},
};
struct ggml_opencl_fa_dim_table {
File diff suppressed because it is too large Load Diff
@@ -10,7 +10,12 @@
#define DK_VEC (DK/4)
#define DV_VEC (DV/4)
#define WG_SIZE (BLOCK_M)
#define Q1_WG_SIZE 64
// q1 reduces over a Q1_WG_SIZE-wide WG via work-group barriers; the launch WG
// must match. Defaults to the Adreno sg (64); host passes -D FA_SG=32 on Intel.
#ifndef FA_SG
#define FA_SG 64
#endif
#define Q1_WG_SIZE FA_SG
// The kernels are built with -cl-finite-math-only. On some older Adreno GPUs,
// infinite operand can cause undefined behavior and miscompilation for exp.
+15 -1
View File
@@ -11,7 +11,12 @@
#define DK_VEC (DK/4)
#define DV_VEC (DV/4)
#define WG_SIZE (BLOCK_M)
#define Q1_WG_SIZE 64
// q1 reduces over a Q1_WG_SIZE-wide WG via work-group barriers; the launch WG
// must match. Defaults to the Adreno sg (64); host passes -D FA_SG=32 on Intel.
#ifndef FA_SG
#define FA_SG 64
#endif
#define Q1_WG_SIZE FA_SG
// The kernels are built with -cl-finite-math-only. On some older Adreno GPUs,
// infinite operand can cause undefined behavior and miscompilation for exp.
@@ -114,6 +119,15 @@ __kernel void flash_attn_f32(
__local DATA_TYPE4 l_v[BLOCK_N][DV_VEC];
for (int k_start = 0; k_start < n_kv; k_start += BLOCK_N) {
#if FA_SG < 64
// WAR on l_k/l_v: threads with my_query_row >= n_q skip the compute below
// (continue) and would race ahead to reload the tiles while active threads
// still read them. A single 64-wide Adreno subgroup (WG == sg) runs lockstep
// and hides this; a WG that spans multiple narrower subgroups (Intel sg=32)
// corrupts the result. All threads reach this each iteration (no-op on the
// first), so it does not diverge with the continue. Compiled out at sg=64.
barrier(CLK_LOCAL_MEM_FENCE);
#endif
for (int i = tid; i < BLOCK_N * DK_VEC; i += WG_SIZE) {
const int row = i / DK_VEC;
const int col = i % DK_VEC;
File diff suppressed because it is too large Load Diff
@@ -27,7 +27,11 @@
#define DK_VEC (DK/4)
#define DV_VEC (DV/4)
#define Q1_WG_SIZE 64
#ifndef FA_SG
#define FA_SG 64
#endif
#define Q1_WG_SIZE FA_SG
// The kernels are built with -cl-finite-math-only. On some older Adreno GPUs,
// infinite operand can cause undefined behavior and miscompilation for exp.
@@ -365,6 +369,263 @@ __kernel void flash_attn_f32_q4_0_q1(
}
}
#ifdef cl_intel_subgroups
#pragma OPENCL EXTENSION cl_intel_subgroups : enable
#else
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#endif
#ifdef cl_qcom_reqd_sub_group_size
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
#else
#define REQD_SUBGROUP_SIZE_64
#endif
#define VEC_NSG 4
#define VEC_WG_SIZE (Q1_WG_SIZE * VEC_NSG)
#define Q1V_DV_PER_THREAD ((DV_VEC + Q1_WG_SIZE - 1) / Q1_WG_SIZE)
// Dequant one float4 lane (0..7) from a q4_0 block.
// Lanes 0..3 low nibbles of qs[0..15], lanes 4..7 high nibbles.
inline float4 dequant_q4_0_lane(const global char * block_ptr, int lane) {
const float d = vload_half(0, (const global half *)block_ptr);
const global uchar * qs = (const global uchar *)(block_ptr + 2);
const int g = lane & 3;
const int shift = (lane < 4) ? 0 : 4;
return d * (float4)((float)((qs[g*4+0] >> shift) & 0x0F) - 8.0f,
(float)((qs[g*4+1] >> shift) & 0x0F) - 8.0f,
(float)((qs[g*4+2] >> shift) & 0x0F) - 8.0f,
(float)((qs[g*4+3] >> shift) & 0x0F) - 8.0f);
}
REQD_SUBGROUP_SIZE_64
__kernel void flash_attn_f32_q4_0_q1_vec(
const global void * q_void, ulong q_offset,
const global void * k_void, ulong k_offset,
const global void * v_void, ulong v_offset,
global void * o_void, ulong o_offset,
const float scale,
const int n_q,
const int n_kv,
const int is_causal,
const int n_head,
const ulong q_nb1, const ulong q_nb2, const ulong q_nb3,
const ulong k_nb1, const ulong k_nb2, const ulong k_nb3,
const ulong v_nb1, const ulong v_nb2, const ulong v_nb3,
const ulong o_nb1, const ulong o_nb2, const ulong o_nb3,
const float max_bias,
const float m0,
const float m1,
const int n_head_log2,
const float logit_softcap,
const int n_head_kv,
const global void* mask_void,
const ulong mask_offset,
const ulong mask_nb1,
const ulong mask_nb2,
const ulong mask_nb3,
const int mask_ne2,
const int mask_ne3,
const global void* sinks_void,
const ulong sinks_offset
) {
const int tid = get_local_id(0);
const int sgid = tid / Q1_WG_SIZE;
const int tid_sg = tid % Q1_WG_SIZE;
const int head_batch_idx = get_global_id(1);
const int batch_idx = head_batch_idx / n_head;
const int head_idx = head_batch_idx % n_head;
const int gqa_ratio = n_head / n_head_kv;
const int head_kv_idx = head_idx / gqa_ratio;
const global char * q_base = (const global char *) q_void + q_offset;
const global char * k_base = (const global char *) k_void + k_offset;
const global char * v_base = (const global char *) v_void + v_offset;
global char * o_base = (global char *) o_void + o_offset;
const global char * mask_base = NULL;
if (mask_void != NULL) {
const int mask_head_idx = head_idx % mask_ne2;
const int mask_batch_idx = batch_idx % mask_ne3;
mask_base = (const global char *) mask_void + mask_offset +
mask_batch_idx * mask_nb3 + mask_head_idx * mask_nb2;
}
__local ACC_TYPE4 q_shared[DK_VEC];
{
const ulong q_row_offset = batch_idx * q_nb3 + head_idx * q_nb2;
const global Q_DATA_TYPE4 * q_ptr = (const global Q_DATA_TYPE4 *) (q_base + q_row_offset);
for (int i = tid; i < DK_VEC; i += VEC_WG_SIZE) {
q_shared[i] = CONVERT_Q_ACC4(q_ptr[i]);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
#ifdef FA_HAVE_INT_DOT
// quantize Q to int8-packed uints + per-block (qd, q_sum) once per WG for dp4a
// one thread per Q block, remaining threads idle this step
__local uint q_packed_shared[DK_Q4_BLOCKS * 8];
__local float q_d_shared[DK_Q4_BLOCKS];
__local int q_sum_shared[DK_Q4_BLOCKS];
if (tid < DK_Q4_BLOCKS) {
ACC_TYPE4 q_block[8];
#pragma unroll
for (int i = 0; i < 8; ++i) q_block[i] = q_shared[tid * 8 + i];
uint packed[8];
q4_q_block_info info = quant_q_block_int8_packed_q4(q_block, packed);
#pragma unroll
for (int i = 0; i < 8; ++i) q_packed_shared[tid * 8 + i] = packed[i];
q_d_shared[tid] = info.qd;
q_sum_shared[tid] = info.q_sum;
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
const float slope = get_alibi_slope(max_bias, head_idx, n_head_log2, m0, m1);
const global ACC_TYPE * sinks_ptr = NULL;
if (sinks_void != NULL) {
sinks_ptr = (const global ACC_TYPE *) ((const global char *) sinks_void + sinks_offset);
}
ACC_TYPE4 o_acc[Q1V_DV_PER_THREAD];
#pragma unroll
for (int i = 0; i < Q1V_DV_PER_THREAD; ++i) o_acc[i] = (ACC_TYPE4)(0.0f);
ACC_TYPE m_i = FA_M_INIT;
ACC_TYPE l_i = 0.0f;
const int kv_per_sg = (n_kv + VEC_NSG - 1) / VEC_NSG;
const int kv_start = sgid * kv_per_sg;
const int kv_end = min(n_kv, kv_start + kv_per_sg);
for (int k_idx = kv_start; k_idx < kv_end; ++k_idx) {
const global char * k_row = k_base + batch_idx * k_nb3 + head_kv_idx * k_nb2 + k_idx * k_nb1;
const global char * v_row = v_base + batch_idx * v_nb3 + head_kv_idx * v_nb2 + k_idx * v_nb1;
#ifdef FA_HAVE_INT_DOT
// per-lane dp4a: each lane packs 4 raw q4_0 nibbles into a uint,
// then dot_acc_sat_4x8packed_ss_int against the matching uint.
ACC_TYPE lane_contrib = 0.0f;
for (int qk = tid_sg; qk < DK_VEC; qk += Q1_WG_SIZE) {
const int block_idx = qk / 8;
const int lane_in_block = qk % 8;
const int g = lane_in_block & 3;
const int shift = (lane_in_block < 4) ? 0 : 4;
const global char * k_block = k_row + block_idx * Q4_0_BLOCK_SIZE;
const float kd = vload_half(0, (const global half *)k_block);
const global uchar * k_qs = (const global uchar *)(k_block + 2);
const uchar b0 = k_qs[g*4 + 0];
const uchar b1 = k_qs[g*4 + 1];
const uchar b2 = k_qs[g*4 + 2];
const uchar b3 = k_qs[g*4 + 3];
const uint k_packed = ((uint)((b0 >> shift) & 0x0F)) |
((uint)((b1 >> shift) & 0x0F)) << 8 |
((uint)((b2 >> shift) & 0x0F)) << 16 |
((uint)((b3 >> shift) & 0x0F)) << 24;
const uint q_packed_lane = q_packed_shared[block_idx * 8 + lane_in_block];
const int raw_dot = dot_acc_sat_4x8packed_ss_int(q_packed_lane, k_packed, 0);
const float qd = q_d_shared[block_idx];
const float block_scale = qd * kd;
float contrib = (float)raw_dot * block_scale;
if (lane_in_block == 0) {
// block bias correction is per-block
const int q_sum_b = q_sum_shared[block_idx];
contrib -= 8.0f * block_scale * (float)q_sum_b;
}
lane_contrib += contrib;
}
ACC_TYPE score = sub_group_reduce_add(lane_contrib) * scale;
#else
ACC_TYPE4 dot4 = (ACC_TYPE4)(0.0f);
for (int qk = tid_sg; qk < DK_VEC; qk += Q1_WG_SIZE) {
const int block_idx = qk / 8;
const int lane = qk % 8;
const float4 k_v = dequant_q4_0_lane(k_row + block_idx * Q4_0_BLOCK_SIZE, lane);
dot4 = mad(q_shared[qk], k_v, dot4);
}
ACC_TYPE dot_partial = dot4.s0 + dot4.s1 + dot4.s2 + dot4.s3;
ACC_TYPE score = sub_group_reduce_add(dot_partial) * scale;
#endif
if (mask_base != NULL) {
const global MASK_DATA_TYPE * mask_ptr = (const global MASK_DATA_TYPE *) mask_base;
score += slope * (ACC_TYPE) mask_ptr[k_idx];
}
if (logit_softcap > 0.0f) {
score = logit_softcap * tanh(score / logit_softcap);
}
const ACC_TYPE m_new = max(m_i, score);
const ACC_TYPE scale_prev = native_exp(m_i - m_new);
const ACC_TYPE p = native_exp(score - m_new);
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
const int block_idx = dv / 8;
const int lane = dv % 8;
const float4 v_v = dequant_q4_0_lane(v_row + block_idx * Q4_0_BLOCK_SIZE, lane);
o_acc[idx] = mad(p, v_v, o_acc[idx] * scale_prev);
}
l_i = l_i * scale_prev + p;
m_i = m_new;
}
__local ACC_TYPE sg_m[VEC_NSG];
__local ACC_TYPE sg_l[VEC_NSG];
__local ACC_TYPE4 sg_o[VEC_NSG][DV_VEC];
if (tid_sg == 0) {
sg_m[sgid] = m_i;
sg_l[sgid] = l_i;
}
{
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
sg_o[sgid][dv] = o_acc[idx];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if (sgid == 0) {
ACC_TYPE m_final = sg_m[0];
#pragma unroll
for (int s = 1; s < VEC_NSG; ++s) {
m_final = max(m_final, sg_m[s]);
}
if (sinks_ptr != NULL) {
m_final = max(m_final, sinks_ptr[head_idx]);
}
ACC_TYPE l_final = 0.0f;
#pragma unroll
for (int s = 0; s < VEC_NSG; ++s) {
l_final += sg_l[s] * native_exp(sg_m[s] - m_final);
}
if (sinks_ptr != NULL) {
l_final += native_exp(sinks_ptr[head_idx] - m_final);
}
const ACC_TYPE l_inv = (l_final > 0.0f) ? (1.0f / l_final) : 0.0f;
const ulong o_row_offset = batch_idx * o_nb3 + head_idx * o_nb1;
global O_DATA_TYPE4 * o_row = (global O_DATA_TYPE4 *) (o_base + o_row_offset);
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
ACC_TYPE4 o_merged = (ACC_TYPE4)(0.0f);
#pragma unroll
for (int s = 0; s < VEC_NSG; ++s) {
const ACC_TYPE alpha = native_exp(sg_m[s] - m_final);
o_merged = mad((ACC_TYPE4)(alpha), sg_o[s][dv], o_merged);
}
o_row[dv] = CONVERT_O_DATA4(o_merged * l_inv);
}
}
}
// Flash-decoding split pass for q4_0 KV. Merge kernel is type-agnostic and
// shared with the f16/q8_0 FA kernels.
#define FA_PARTIAL_FLOATS (2 + DV)
@@ -583,6 +844,319 @@ __kernel void flash_attn_f32_q4_0_q1_split(
#define WG_SIZE BLOCK_M
#endif
#ifndef MQ_GQA
#define MQ_GQA 4
#endif
#ifndef MQ_NSG_SPLIT
#define MQ_NSG_SPLIT 4
#endif
#define MQ_SPLIT_WG_SIZE_Q4 (Q1_WG_SIZE * MQ_NSG_SPLIT)
REQD_SUBGROUP_SIZE_64
__kernel void flash_attn_f32_q4_0_q1_vec_mq_split(
const global void * q_void, ulong q_offset,
const global void * k_void, ulong k_offset,
const global void * v_void, ulong v_offset,
const float scale,
const int n_q,
const int n_kv,
const int n_head,
const ulong q_nb1, const ulong q_nb2, const ulong q_nb3,
const ulong k_nb1, const ulong k_nb2, const ulong k_nb3,
const ulong v_nb1, const ulong v_nb2, const ulong v_nb3,
const float max_bias,
const float m0,
const float m1,
const int n_head_log2,
const float logit_softcap,
const int n_head_kv,
const global void * mask_void,
const ulong mask_offset,
const ulong mask_nb1,
const ulong mask_nb2,
const ulong mask_nb3,
const int mask_ne2,
const int mask_ne3,
global float * partial_void,
const int n_splits,
const int kv_per_split
) {
const int tid = get_local_id(0);
const int sgid = tid / Q1_WG_SIZE;
const int tid_sg = tid % Q1_WG_SIZE;
const int kvhead_batch_idx = get_global_id(1);
const int split_q_idx = get_global_id(2);
const int split_idx = split_q_idx % n_splits;
const int q_idx = split_q_idx / n_splits;
const int batch_idx = kvhead_batch_idx / n_head_kv;
const int head_kv_idx = kvhead_batch_idx % n_head_kv;
const int kv_start = split_idx * kv_per_split;
const int kv_end = min(kv_start + kv_per_split, n_kv);
const ulong record_stride = (ulong) FA_PARTIAL_FLOATS;
if (kv_start >= kv_end) {
if (tid == 0) {
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const int head_idx = head_kv_idx * MQ_GQA + h;
const ulong rec_idx = ((((ulong) batch_idx * n_head + head_idx) * n_q + q_idx)
* n_splits + split_idx);
global float * rec = partial_void + rec_idx * record_stride;
rec[0] = FA_M_INIT;
rec[1] = 0.0f;
}
}
return;
}
const global char * q_base = (const global char *) q_void + q_offset;
const global char * k_base = (const global char *) k_void + k_offset;
const global char * v_base = (const global char *) v_void + v_offset;
__local ACC_TYPE4 q_shared[MQ_GQA * DK_VEC];
for (int i = tid; i < MQ_GQA * DK_VEC; i += MQ_SPLIT_WG_SIZE_Q4) {
const int h = i / DK_VEC;
const int k = i % DK_VEC;
const int head_idx = head_kv_idx * MQ_GQA + h;
const ulong q_row_offset = batch_idx * q_nb3 + head_idx * q_nb2 + (ulong) q_idx * q_nb1;
const global Q_DATA_TYPE4 * q_ptr = (const global Q_DATA_TYPE4 *) (q_base + q_row_offset);
q_shared[h * DK_VEC + k] = CONVERT_Q_ACC4(q_ptr[k]);
}
barrier(CLK_LOCAL_MEM_FENCE);
#ifdef FA_HAVE_INT_DOT
__local uint q_packed_shared[MQ_GQA * DK_Q4_BLOCKS * 8];
__local float q_d_shared[MQ_GQA * DK_Q4_BLOCKS];
__local int q_sum_shared[MQ_GQA * DK_Q4_BLOCKS];
{
const int active = MQ_GQA * DK_Q4_BLOCKS;
if (tid < active) {
const int h = tid / DK_Q4_BLOCKS;
const int block_id = tid % DK_Q4_BLOCKS;
ACC_TYPE4 q_block[8];
#pragma unroll
for (int i = 0; i < 8; ++i) q_block[i] = q_shared[h * DK_VEC + block_id * 8 + i];
uint packed[8];
q4_q_block_info info = quant_q_block_int8_packed_q4(q_block, packed);
#pragma unroll
for (int i = 0; i < 8; ++i) q_packed_shared[(h * DK_Q4_BLOCKS + block_id) * 8 + i] = packed[i];
q_d_shared[h * DK_Q4_BLOCKS + block_id] = info.qd;
q_sum_shared[h * DK_Q4_BLOCKS + block_id] = info.q_sum;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
float slope[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
slope[h] = get_alibi_slope(max_bias, head_kv_idx * MQ_GQA + h, n_head_log2, m0, m1);
}
const global char * mask_base[MQ_GQA];
if (mask_void != NULL) {
const int mask_batch_idx = batch_idx % mask_ne3;
const global char * mask_base_b = (const global char *) mask_void + mask_offset +
mask_batch_idx * mask_nb3 +
(ulong) q_idx * mask_nb1;
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const int head_idx = head_kv_idx * MQ_GQA + h;
const int mask_head_idx = head_idx % mask_ne2;
mask_base[h] = mask_base_b + mask_head_idx * mask_nb2;
}
} else {
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) mask_base[h] = NULL;
}
ACC_TYPE4 o_acc[MQ_GQA][Q1V_DV_PER_THREAD];
ACC_TYPE m_i[MQ_GQA];
ACC_TYPE l_i[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
m_i[h] = FA_M_INIT;
l_i[h] = 0.0f;
#pragma unroll
for (int i = 0; i < Q1V_DV_PER_THREAD; ++i) o_acc[h][i] = (ACC_TYPE4)(0.0f);
}
const int kv_len = kv_end - kv_start;
const int kv_per_sg = (kv_len + MQ_NSG_SPLIT - 1) / MQ_NSG_SPLIT;
const int kv_lo = kv_start + sgid * kv_per_sg;
const int kv_hi = min(kv_end, kv_lo + kv_per_sg);
for (int k_idx = kv_lo; k_idx < kv_hi; ++k_idx) {
const global char * k_row = k_base + batch_idx * k_nb3 + head_kv_idx * k_nb2 + k_idx * k_nb1;
const global char * v_row = v_base + batch_idx * v_nb3 + head_kv_idx * v_nb2 + k_idx * v_nb1;
#ifdef FA_HAVE_INT_DOT
ACC_TYPE lane_contrib[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) lane_contrib[h] = 0.0f;
for (int qk = tid_sg; qk < DK_VEC; qk += Q1_WG_SIZE) {
const int block_idx = qk / 8;
const int lane_in_block = qk % 8;
const int g = lane_in_block & 3;
const int shift = (lane_in_block < 4) ? 0 : 4;
const global char * k_block = k_row + block_idx * Q4_0_BLOCK_SIZE;
const float kd = vload_half(0, (const global half *)k_block);
const global uchar * k_qs = (const global uchar *)(k_block + 2);
const uchar b0 = k_qs[g*4 + 0];
const uchar b1 = k_qs[g*4 + 1];
const uchar b2 = k_qs[g*4 + 2];
const uchar b3 = k_qs[g*4 + 3];
const uint k_packed = ((uint)((b0 >> shift) & 0x0F)) |
((uint)((b1 >> shift) & 0x0F)) << 8 |
((uint)((b2 >> shift) & 0x0F)) << 16 |
((uint)((b3 >> shift) & 0x0F)) << 24;
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const uint q_packed_lane = q_packed_shared[(h * DK_Q4_BLOCKS + block_idx) * 8 + lane_in_block];
const int raw_dot = dot_acc_sat_4x8packed_ss_int(q_packed_lane, k_packed, 0);
const float qd = q_d_shared[h * DK_Q4_BLOCKS + block_idx];
const float block_scale = qd * kd;
float contrib = (float) raw_dot * block_scale;
if (lane_in_block == 0) {
const int q_sum_b = q_sum_shared[h * DK_Q4_BLOCKS + block_idx];
contrib -= 8.0f * block_scale * (float) q_sum_b;
}
lane_contrib[h] += contrib;
}
}
ACC_TYPE score[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
ACC_TYPE s = sub_group_reduce_add(lane_contrib[h]) * scale;
if (mask_base[h] != NULL) {
const global MASK_DATA_TYPE * mask_ptr = (const global MASK_DATA_TYPE *) mask_base[h];
s += slope[h] * (ACC_TYPE) mask_ptr[k_idx];
}
if (logit_softcap > 0.0f) {
s = logit_softcap * tanh(s / logit_softcap);
}
score[h] = s;
}
#else
// fallback float-dequant K dot
ACC_TYPE4 dot4[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) dot4[h] = (ACC_TYPE4)(0.0f);
for (int qk = tid_sg; qk < DK_VEC; qk += Q1_WG_SIZE) {
const int block_idx = qk / 8;
const int lane = qk % 8;
const float4 k_v = dequant_q4_0_lane(k_row + block_idx * Q4_0_BLOCK_SIZE, lane);
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
dot4[h] = mad(q_shared[h * DK_VEC + qk], k_v, dot4[h]);
}
}
ACC_TYPE score[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const ACC_TYPE dot_partial = dot4[h].s0 + dot4[h].s1 + dot4[h].s2 + dot4[h].s3;
ACC_TYPE s = sub_group_reduce_add(dot_partial) * scale;
if (mask_base[h] != NULL) {
const global MASK_DATA_TYPE * mask_ptr = (const global MASK_DATA_TYPE *) mask_base[h];
s += slope[h] * (ACC_TYPE) mask_ptr[k_idx];
}
if (logit_softcap > 0.0f) {
s = logit_softcap * tanh(s / logit_softcap);
}
score[h] = s;
}
#endif
ACC_TYPE p_h[MQ_GQA];
ACC_TYPE sp_h[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const ACC_TYPE m_new = max(m_i[h], score[h]);
sp_h[h] = native_exp(m_i[h] - m_new);
p_h[h] = native_exp(score[h] - m_new);
l_i[h] = l_i[h] * sp_h[h] + p_h[h];
m_i[h] = m_new;
}
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
const int block_idx = dv / 8;
const int lane = dv % 8;
const float4 v_v = dequant_q4_0_lane(v_row + block_idx * Q4_0_BLOCK_SIZE, lane);
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
o_acc[h][idx] = mad(p_h[h], v_v, o_acc[h][idx] * sp_h[h]);
}
}
}
// per-h cross-subgroup merge
__local ACC_TYPE sg_m[MQ_GQA][MQ_NSG_SPLIT];
__local ACC_TYPE sg_l[MQ_GQA][MQ_NSG_SPLIT];
__local ACC_TYPE4 sg_o[MQ_NSG_SPLIT][DV_VEC];
if (tid_sg == 0) {
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
sg_m[h][sgid] = m_i[h];
sg_l[h][sgid] = l_i[h];
}
}
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
{
int idx = 0;
for (int dv_idx = tid_sg; dv_idx < DV_VEC; dv_idx += Q1_WG_SIZE, ++idx) {
sg_o[sgid][dv_idx] = o_acc[h][idx];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if (sgid == 0) {
const int head_idx = head_kv_idx * MQ_GQA + h;
ACC_TYPE m_c = sg_m[h][0];
#pragma unroll
for (int s = 1; s < MQ_NSG_SPLIT; ++s) {
m_c = max(m_c, sg_m[h][s]);
}
ACC_TYPE l_c = 0.0f;
#pragma unroll
for (int s = 0; s < MQ_NSG_SPLIT; ++s) {
l_c += sg_l[h][s] * native_exp(sg_m[h][s] - m_c);
}
const ulong rec_idx = ((((ulong) batch_idx * n_head + head_idx) * n_q + q_idx)
* n_splits + split_idx);
global float * rec = partial_void + rec_idx * record_stride;
global float4 * rec_o = (global float4 *) (rec + 2);
if (tid_sg == 0) {
rec[0] = (float) m_c;
rec[1] = (float) l_c;
}
for (int dv_idx = tid_sg; dv_idx < DV_VEC; dv_idx += Q1_WG_SIZE) {
ACC_TYPE4 o_merged = (ACC_TYPE4)(0.0f);
#pragma unroll
for (int s = 0; s < MQ_NSG_SPLIT; ++s) {
const ACC_TYPE alpha = native_exp(sg_m[h][s] - m_c);
o_merged = mad((ACC_TYPE4)(alpha), sg_o[s][dv_idx], o_merged);
}
rec_o[dv_idx] = o_merged;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}
__kernel void flash_attn_f32_q4_0(
const global void * q_void, ulong q_offset,
const global void * k_void, ulong k_offset,
@@ -24,7 +24,11 @@
#define DK_VEC (DK/4)
#define DV_VEC (DV/4)
#define Q1_WG_SIZE 64
#ifndef FA_SG
#define FA_SG 64
#endif
#define Q1_WG_SIZE FA_SG
// The kernels are built with -cl-finite-math-only. On some older Adreno GPUs,
// infinite operand can cause undefined behavior and miscompilation for exp.
@@ -310,6 +314,201 @@ __kernel void flash_attn_f32_q8_0_q1(
}
}
#ifdef cl_intel_subgroups
#pragma OPENCL EXTENSION cl_intel_subgroups : enable
#else
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#endif
#ifdef cl_qcom_reqd_sub_group_size
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
#else
#define REQD_SUBGROUP_SIZE_64
#endif
#define VEC_NSG 4
#define VEC_WG_SIZE (Q1_WG_SIZE * VEC_NSG)
#define Q1V_DV_PER_THREAD ((DV_VEC + Q1_WG_SIZE - 1) / Q1_WG_SIZE)
inline float4 dequant_q8_0_lane(const global char * block_ptr, int lane) {
const float d = vload_half(0, (const global half *)block_ptr);
const global char * qs = block_ptr + 2 + lane * 4;
return d * (float4)((float)qs[0], (float)qs[1], (float)qs[2], (float)qs[3]);
}
REQD_SUBGROUP_SIZE_64
__kernel void flash_attn_f32_q8_0_q1_vec(
const global void * q_void, ulong q_offset,
const global void * k_void, ulong k_offset,
const global void * v_void, ulong v_offset,
global void * o_void, ulong o_offset,
const float scale,
const int n_q,
const int n_kv,
const int is_causal,
const int n_head,
const ulong q_nb1, const ulong q_nb2, const ulong q_nb3,
const ulong k_nb1, const ulong k_nb2, const ulong k_nb3,
const ulong v_nb1, const ulong v_nb2, const ulong v_nb3,
const ulong o_nb1, const ulong o_nb2, const ulong o_nb3,
const float max_bias,
const float m0,
const float m1,
const int n_head_log2,
const float logit_softcap,
const int n_head_kv,
const global void* mask_void,
const ulong mask_offset,
const ulong mask_nb1,
const ulong mask_nb2,
const ulong mask_nb3,
const int mask_ne2,
const int mask_ne3,
const global void* sinks_void,
const ulong sinks_offset
) {
const int tid = get_local_id(0);
const int sgid = tid / Q1_WG_SIZE;
const int tid_sg = tid % Q1_WG_SIZE;
const int head_batch_idx = get_global_id(1);
const int batch_idx = head_batch_idx / n_head;
const int head_idx = head_batch_idx % n_head;
const int gqa_ratio = n_head / n_head_kv;
const int head_kv_idx = head_idx / gqa_ratio;
const global char * q_base = (const global char *) q_void + q_offset;
const global char * k_base = (const global char *) k_void + k_offset;
const global char * v_base = (const global char *) v_void + v_offset;
global char * o_base = (global char *) o_void + o_offset;
const global char * mask_base = NULL;
if (mask_void != NULL) {
const int mask_head_idx = head_idx % mask_ne2;
const int mask_batch_idx = batch_idx % mask_ne3;
mask_base = (const global char *) mask_void + mask_offset +
mask_batch_idx * mask_nb3 + mask_head_idx * mask_nb2;
}
__local ACC_TYPE4 q_shared[DK_VEC];
{
const ulong q_row_offset = batch_idx * q_nb3 + head_idx * q_nb2;
const global Q_DATA_TYPE4 * q_ptr = (const global Q_DATA_TYPE4 *) (q_base + q_row_offset);
for (int i = tid; i < DK_VEC; i += VEC_WG_SIZE) {
q_shared[i] = CONVERT_Q_ACC4(q_ptr[i]);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
const float slope = get_alibi_slope(max_bias, head_idx, n_head_log2, m0, m1);
const global ACC_TYPE * sinks_ptr = NULL;
if (sinks_void != NULL) {
sinks_ptr = (const global ACC_TYPE *) ((const global char *) sinks_void + sinks_offset);
}
ACC_TYPE4 o_acc[Q1V_DV_PER_THREAD];
#pragma unroll
for (int i = 0; i < Q1V_DV_PER_THREAD; ++i) o_acc[i] = (ACC_TYPE4)(0.0f);
ACC_TYPE m_i = FA_M_INIT;
ACC_TYPE l_i = 0.0f;
const int kv_per_sg = (n_kv + VEC_NSG - 1) / VEC_NSG;
const int kv_start = sgid * kv_per_sg;
const int kv_end = min(n_kv, kv_start + kv_per_sg);
for (int k_idx = kv_start; k_idx < kv_end; ++k_idx) {
const global char * k_row = k_base + batch_idx * k_nb3 + head_kv_idx * k_nb2 + k_idx * k_nb1;
const global char * v_row = v_base + batch_idx * v_nb3 + head_kv_idx * v_nb2 + k_idx * v_nb1;
ACC_TYPE4 dot4 = (ACC_TYPE4)(0.0f);
for (int qk = tid_sg; qk < DK_VEC; qk += Q1_WG_SIZE) {
const int block_idx = qk / 8;
const int lane = qk % 8;
const float4 k_v = dequant_q8_0_lane(k_row + block_idx * Q8_0_BLOCK_SIZE, lane);
dot4 = mad(q_shared[qk], k_v, dot4);
}
ACC_TYPE dot_partial = dot4.s0 + dot4.s1 + dot4.s2 + dot4.s3;
ACC_TYPE score = sub_group_reduce_add(dot_partial) * scale;
if (mask_base != NULL) {
const global MASK_DATA_TYPE * mask_ptr = (const global MASK_DATA_TYPE *) mask_base;
score += slope * (ACC_TYPE) mask_ptr[k_idx];
}
if (logit_softcap > 0.0f) {
score = logit_softcap * tanh(score / logit_softcap);
}
const ACC_TYPE m_new = max(m_i, score);
const ACC_TYPE scale_prev = native_exp(m_i - m_new);
const ACC_TYPE p = native_exp(score - m_new);
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
const int block_idx = dv / 8;
const int lane = dv % 8;
const float4 v_v = dequant_q8_0_lane(v_row + block_idx * Q8_0_BLOCK_SIZE, lane);
o_acc[idx] = mad(p, v_v, o_acc[idx] * scale_prev);
}
l_i = l_i * scale_prev + p;
m_i = m_new;
}
__local ACC_TYPE sg_m[VEC_NSG];
__local ACC_TYPE sg_l[VEC_NSG];
__local ACC_TYPE4 sg_o[VEC_NSG][DV_VEC];
if (tid_sg == 0) {
sg_m[sgid] = m_i;
sg_l[sgid] = l_i;
}
{
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
sg_o[sgid][dv] = o_acc[idx];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if (sgid == 0) {
ACC_TYPE m_final = sg_m[0];
#pragma unroll
for (int s = 1; s < VEC_NSG; ++s) {
m_final = max(m_final, sg_m[s]);
}
if (sinks_ptr != NULL) {
m_final = max(m_final, sinks_ptr[head_idx]);
}
ACC_TYPE l_final = 0.0f;
#pragma unroll
for (int s = 0; s < VEC_NSG; ++s) {
l_final += sg_l[s] * native_exp(sg_m[s] - m_final);
}
if (sinks_ptr != NULL) {
l_final += native_exp(sinks_ptr[head_idx] - m_final);
}
const ACC_TYPE l_inv = (l_final > 0.0f) ? (1.0f / l_final) : 0.0f;
const ulong o_row_offset = batch_idx * o_nb3 + head_idx * o_nb1;
global O_DATA_TYPE4 * o_row = (global O_DATA_TYPE4 *) (o_base + o_row_offset);
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
ACC_TYPE4 o_merged = (ACC_TYPE4)(0.0f);
#pragma unroll
for (int s = 0; s < VEC_NSG; ++s) {
const ACC_TYPE alpha = native_exp(sg_m[s] - m_final);
o_merged = mad((ACC_TYPE4)(alpha), sg_o[s][dv], o_merged);
}
o_row[dv] = CONVERT_O_DATA4(o_merged * l_inv);
}
}
}
// Flash-decoding split pass for q8_0 KV. Partial record: [m, l, O[DV]].
// Merge kernel from flash_attn_f32_f16.cl is type-agnostic and reused.
#define FA_PARTIAL_FLOATS (2 + DV)
@@ -533,6 +732,244 @@ __kernel void flash_attn_f32_q8_0_q1_split(
#define FA_V_STRATEGY 0
#endif
#ifndef MQ_GQA
#define MQ_GQA 4
#endif
#ifndef MQ_NSG_SPLIT
#define MQ_NSG_SPLIT 4
#endif
#define MQ_SPLIT_WG_SIZE_Q8 (Q1_WG_SIZE * MQ_NSG_SPLIT)
REQD_SUBGROUP_SIZE_64
__kernel void flash_attn_f32_q8_0_q1_vec_mq_split(
const global void * q_void, ulong q_offset,
const global void * k_void, ulong k_offset,
const global void * v_void, ulong v_offset,
const float scale,
const int n_q,
const int n_kv,
const int n_head,
const ulong q_nb1, const ulong q_nb2, const ulong q_nb3,
const ulong k_nb1, const ulong k_nb2, const ulong k_nb3,
const ulong v_nb1, const ulong v_nb2, const ulong v_nb3,
const float max_bias,
const float m0,
const float m1,
const int n_head_log2,
const float logit_softcap,
const int n_head_kv,
const global void * mask_void,
const ulong mask_offset,
const ulong mask_nb1,
const ulong mask_nb2,
const ulong mask_nb3,
const int mask_ne2,
const int mask_ne3,
global float * partial_void,
const int n_splits,
const int kv_per_split
) {
const int tid = get_local_id(0);
const int sgid = tid / Q1_WG_SIZE;
const int tid_sg = tid % Q1_WG_SIZE;
const int kvhead_batch_idx = get_global_id(1);
const int split_q_idx = get_global_id(2);
const int split_idx = split_q_idx % n_splits;
const int q_idx = split_q_idx / n_splits;
const int batch_idx = kvhead_batch_idx / n_head_kv;
const int head_kv_idx = kvhead_batch_idx % n_head_kv;
const int kv_start = split_idx * kv_per_split;
const int kv_end = min(kv_start + kv_per_split, n_kv);
const ulong record_stride = (ulong) FA_PARTIAL_FLOATS;
if (kv_start >= kv_end) {
// Empty split write sentinel for each of the MQ_GQA Q-heads.
if (tid == 0) {
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const int head_idx = head_kv_idx * MQ_GQA + h;
const ulong rec_idx = ((((ulong) batch_idx * n_head + head_idx) * n_q + q_idx)
* n_splits + split_idx);
global float * rec = partial_void + rec_idx * record_stride;
rec[0] = FA_M_INIT;
rec[1] = 0.0f;
}
}
return;
}
const global char * q_base = (const global char *) q_void + q_offset;
const global char * k_base = (const global char *) k_void + k_offset;
const global char * v_base = (const global char *) v_void + v_offset;
__local ACC_TYPE4 q_shared[MQ_GQA * DK_VEC];
for (int i = tid; i < MQ_GQA * DK_VEC; i += MQ_SPLIT_WG_SIZE_Q8) {
const int h = i / DK_VEC;
const int k = i % DK_VEC;
const int head_idx = head_kv_idx * MQ_GQA + h;
const ulong q_row_offset = batch_idx * q_nb3 + head_idx * q_nb2 + (ulong) q_idx * q_nb1;
const global Q_DATA_TYPE4 * q_ptr = (const global Q_DATA_TYPE4 *) (q_base + q_row_offset);
q_shared[h * DK_VEC + k] = CONVERT_Q_ACC4(q_ptr[k]);
}
barrier(CLK_LOCAL_MEM_FENCE);
float slope[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
slope[h] = get_alibi_slope(max_bias, head_kv_idx * MQ_GQA + h, n_head_log2, m0, m1);
}
const global char * mask_base[MQ_GQA];
if (mask_void != NULL) {
const int mask_batch_idx = batch_idx % mask_ne3;
const global char * mask_base_b = (const global char *) mask_void + mask_offset +
mask_batch_idx * mask_nb3 +
(ulong) q_idx * mask_nb1;
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const int head_idx = head_kv_idx * MQ_GQA + h;
const int mask_head_idx = head_idx % mask_ne2;
mask_base[h] = mask_base_b + mask_head_idx * mask_nb2;
}
} else {
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) mask_base[h] = NULL;
}
ACC_TYPE4 o_acc[MQ_GQA][Q1V_DV_PER_THREAD];
ACC_TYPE m_i[MQ_GQA];
ACC_TYPE l_i[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
m_i[h] = FA_M_INIT;
l_i[h] = 0.0f;
#pragma unroll
for (int i = 0; i < Q1V_DV_PER_THREAD; ++i) o_acc[h][i] = (ACC_TYPE4)(0.0f);
}
const int kv_len = kv_end - kv_start;
const int kv_per_sg = (kv_len + MQ_NSG_SPLIT - 1) / MQ_NSG_SPLIT;
const int kv_lo = kv_start + sgid * kv_per_sg;
const int kv_hi = min(kv_end, kv_lo + kv_per_sg);
for (int k_idx = kv_lo; k_idx < kv_hi; ++k_idx) {
const global char * k_row = k_base + batch_idx * k_nb3 + head_kv_idx * k_nb2 + k_idx * k_nb1;
const global char * v_row = v_base + batch_idx * v_nb3 + head_kv_idx * v_nb2 + k_idx * v_nb1;
ACC_TYPE4 dot4[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) dot4[h] = (ACC_TYPE4)(0.0f);
for (int qk = tid_sg; qk < DK_VEC; qk += Q1_WG_SIZE) {
const int block_idx = qk / 8;
const int lane = qk % 8;
const float4 k_v = dequant_q8_0_lane(k_row + block_idx * Q8_0_BLOCK_SIZE, lane);
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
dot4[h] = mad(q_shared[h * DK_VEC + qk], k_v, dot4[h]);
}
}
ACC_TYPE score[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const ACC_TYPE dot_partial = dot4[h].s0 + dot4[h].s1 + dot4[h].s2 + dot4[h].s3;
ACC_TYPE s = sub_group_reduce_add(dot_partial) * scale;
if (mask_base[h] != NULL) {
const global MASK_DATA_TYPE * mask_ptr = (const global MASK_DATA_TYPE *) mask_base[h];
s += slope[h] * (ACC_TYPE) mask_ptr[k_idx];
}
if (logit_softcap > 0.0f) {
s = logit_softcap * tanh(s / logit_softcap);
}
score[h] = s;
}
ACC_TYPE p_h[MQ_GQA];
ACC_TYPE sp_h[MQ_GQA];
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
const ACC_TYPE m_new = max(m_i[h], score[h]);
sp_h[h] = native_exp(m_i[h] - m_new);
p_h[h] = native_exp(score[h] - m_new);
l_i[h] = l_i[h] * sp_h[h] + p_h[h];
m_i[h] = m_new;
}
int idx = 0;
for (int dv = tid_sg; dv < DV_VEC; dv += Q1_WG_SIZE, ++idx) {
const int block_idx = dv / 8;
const int lane = dv % 8;
const float4 v_v = dequant_q8_0_lane(v_row + block_idx * Q8_0_BLOCK_SIZE, lane);
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
o_acc[h][idx] = mad(p_h[h], v_v, o_acc[h][idx] * sp_h[h]);
}
}
}
__local ACC_TYPE sg_m[MQ_GQA][MQ_NSG_SPLIT];
__local ACC_TYPE sg_l[MQ_GQA][MQ_NSG_SPLIT];
__local ACC_TYPE4 sg_o[MQ_NSG_SPLIT][DV_VEC];
if (tid_sg == 0) {
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
sg_m[h][sgid] = m_i[h];
sg_l[h][sgid] = l_i[h];
}
}
#pragma unroll
for (int h = 0; h < MQ_GQA; ++h) {
{
int idx = 0;
for (int dv_idx = tid_sg; dv_idx < DV_VEC; dv_idx += Q1_WG_SIZE, ++idx) {
sg_o[sgid][dv_idx] = o_acc[h][idx];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if (sgid == 0) {
const int head_idx = head_kv_idx * MQ_GQA + h;
ACC_TYPE m_c = sg_m[h][0];
#pragma unroll
for (int s = 1; s < MQ_NSG_SPLIT; ++s) {
m_c = max(m_c, sg_m[h][s]);
}
ACC_TYPE l_c = 0.0f;
#pragma unroll
for (int s = 0; s < MQ_NSG_SPLIT; ++s) {
l_c += sg_l[h][s] * native_exp(sg_m[h][s] - m_c);
}
const ulong rec_idx = ((((ulong) batch_idx * n_head + head_idx) * n_q + q_idx)
* n_splits + split_idx);
global float * rec = partial_void + rec_idx * record_stride;
global float4 * rec_o = (global float4 *) (rec + 2);
if (tid_sg == 0) {
rec[0] = (float) m_c;
rec[1] = (float) l_c;
}
for (int dv_idx = tid_sg; dv_idx < DV_VEC; dv_idx += Q1_WG_SIZE) {
ACC_TYPE4 o_merged = (ACC_TYPE4)(0.0f);
#pragma unroll
for (int s = 0; s < MQ_NSG_SPLIT; ++s) {
const ACC_TYPE alpha = native_exp(sg_m[h][s] - m_c);
o_merged = mad((ACC_TYPE4)(alpha), sg_o[s][dv_idx], o_merged);
}
rec_o[dv_idx] = o_merged;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}
__kernel void flash_attn_f32_q8_0(
const global void * q_void, ulong q_offset,
const global void * k_void, ulong k_offset,
@@ -18,6 +18,14 @@
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
#endif
#ifdef cl_khr_subgroup_shuffle
#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable
#define HAS_SUBGROUP_SHUFFLE 1
#elif defined(cl_qcom_subgroup_shuffle)
#pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
#define HAS_SUBGROUP_SHUFFLE 1
#endif
// Assumes row size (ne00) is a multiple of 4
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
@@ -378,3 +386,848 @@ kernel void kernel_mul_mat_f16_f32_l4_dr_lq(
}
}
#endif // ADRENO_GPU
#define N_ROWS_PER_WG 8
#define N_OUTS_PER_WG 8
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_x8(
global char * src0,
ulong offset0,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb00,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src0 = (global char *)((global char *)src0 + offset0);
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int sgs_sz = get_max_sub_group_size();
const int r0_base = get_group_id(0) * N_ROWS_PER_WG;
const int im = get_group_id(2);
const int i12 = im % ne12;
const int i13 = im / ne12;
const ulong offset_src1 = (i12) * nb12 + (i13) * nb13;
global float4 * y4 = (global float4 *)(src1 + offset_src1);
__local float4 q_loc[64]; // ne00/4 max for sub_group_size 64
if (sgs_lid < ne00 / 4) {
q_loc[sgs_lid] = y4[sgs_lid];
}
barrier(CLK_LOCAL_MEM_FENCE);
#pragma unroll
for (int dr = 0; dr < N_ROWS_PER_WG; ++dr) {
const int r0 = r0_base + dr;
if (r0 >= ne01) return;
const ulong offset_src0 = r0 * nb01 + (i12 / r2) * nb02 + (i13 / r3) * nb03;
global half4 * x4 = (global half4 *)(src0 + offset_src0);
float sumf = 0.0f;
for (int i = sgs_lid; i < ne00 / 4; i += sgs_sz) {
const half4 k4 = x4[i];
const float4 q = q_loc[i];
sumf += convert_float(k4.s0) * q.s0
+ convert_float(k4.s1) * q.s1
+ convert_float(k4.s2) * q.s2
+ convert_float(k4.s3) * q.s3;
}
const float all_sum = sub_group_reduce_add(sumf);
if (sgs_lid == 0) {
dst[im * ne1 * ne0 + r0] = all_sum; // ne11 == 1, so r1==0
}
}
}
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_y8(
global char * src0,
ulong offset0,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb00,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src0 = (global char *)((global char *)src0 + offset0);
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int sgs_sz = get_max_sub_group_size();
const int r0_base = get_group_id(0) * N_OUTS_PER_WG;
const int im = get_group_id(2);
const int i12 = im % ne12;
const int i13 = im / ne12;
const ulong offset_src1 = (i12) * nb12 + (i13) * nb13;
global float4 * y4 = (global float4 *)(src1 + offset_src1);
global half4 * x4_o[N_OUTS_PER_WG];
#pragma unroll
for (int o = 0; o < N_OUTS_PER_WG; ++o) {
const int r0 = r0_base + o;
const int r0c = (r0 < ne01) ? r0 : 0;
const ulong off = r0c * nb01 + (i12 / r2) * nb02 + (i13 / r3) * nb03;
x4_o[o] = (global half4 *)(src0 + off);
}
float sum[N_OUTS_PER_WG] = { 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f };
for (int i = sgs_lid; i < ne00 / 4; i += sgs_sz) {
const float4 q4 = y4[i];
#pragma unroll
for (int o = 0; o < N_OUTS_PER_WG; ++o) {
const half4 v4 = x4_o[o][i];
sum[o] += convert_float(v4.s0) * q4.s0
+ convert_float(v4.s1) * q4.s1
+ convert_float(v4.s2) * q4.s2
+ convert_float(v4.s3) * q4.s3;
}
}
#pragma unroll
for (int o = 0; o < N_OUTS_PER_WG; ++o) {
const int r0 = r0_base + o;
const float s = sub_group_reduce_add(sum[o]);
if (sgs_lid == 0 && r0 < ne01) {
dst[im * ne1 * ne0 + r0] = s;
}
}
}
#define N_OUTS_PAIR 8
#define N_PAIRS_PAIR (N_OUTS_PAIR / 2)
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_x8_pair(
global char * src0,
ulong offset0,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb00,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src0 = (global char *)((global char *)src0 + offset0);
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int half_id = sgs_lid >> 5; // 0 = lower half, 1 = upper half
const int lane_h = sgs_lid & 31; // lane 0..31 within half
const int r0_base = get_group_id(0) * N_OUTS_PAIR;
const int im = get_group_id(2);
const int i12 = im % ne12;
const int i13 = im / ne12;
const ulong offset_src1 = (i12) * nb12 + (i13) * nb13;
global float4 * y4 = (global float4 *)(src1 + offset_src1);
__local float4 q_loc[64]; // ne00/4 max for sub_group_size 64
if (sgs_lid < ne00 / 4) {
q_loc[sgs_lid] = y4[sgs_lid];
}
barrier(CLK_LOCAL_MEM_FENCE);
const int dk_vec = ne00 / 4;
#pragma unroll
for (int p = 0; p < N_PAIRS_PAIR; ++p) {
const int r0 = r0_base + 2 * p + half_id;
const ulong offset_src0 = r0 * nb01 + (i12 / r2) * nb02 + (i13 / r3) * nb03;
global half4 * x4 = (global half4 *)(src0 + offset_src0);
float sumf = 0.0f;
for (int i = lane_h; i < dk_vec; i += 32) {
const half4 k4 = x4[i];
const float4 q = q_loc[i];
sumf += convert_float(k4.s0) * q.s0
+ convert_float(k4.s1) * q.s1
+ convert_float(k4.s2) * q.s2
+ convert_float(k4.s3) * q.s3;
}
sumf += sub_group_shuffle_xor(sumf, 16);
sumf += sub_group_shuffle_xor(sumf, 8);
sumf += sub_group_shuffle_xor(sumf, 4);
sumf += sub_group_shuffle_xor(sumf, 2);
sumf += sub_group_shuffle_xor(sumf, 1);
if (lane_h == 0) {
dst[im * ne1 * ne0 + r0] = sumf;
}
}
}
#define N_K_ROWS_GQA 16
#define GQA_RATIO_GQA 8
#define LANES_PER_QH 8 // 64 / GQA_RATIO_GQA
#define DK_VEC_GQA 32 // DK / 4 for DK=128
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_x8_gqa4(
global char * src0,
ulong offset0,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb00,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src0 = (global char *)((global char *)src0 + offset0);
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int q_id = sgs_lid >> 3; // 0..7: which Q-head (8 per WG)
const int lane_q = sgs_lid & 7; // 0..7: lane within Q-head partition
const int r0_base = get_group_id(0) * N_K_ROWS_GQA;
const int im_kv = get_group_id(2);
const int i02 = im_kv % ne02; // K-head index (also K2 batch)
const int i03 = im_kv / ne02; // n13 batch index
const int q_head_lo = i02 * GQA_RATIO_GQA;
__local float4 q_loc[GQA_RATIO_GQA * DK_VEC_GQA]; // 4 × 32 = 128 float4
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_GQA; ++qh) {
const int qh_idx = q_head_lo + qh;
global float4 * y4 = (global float4 *)(src1 + qh_idx * nb12 + i03 * nb13);
if (sgs_lid < DK_VEC_GQA) {
q_loc[qh * DK_VEC_GQA + sgs_lid] = y4[sgs_lid];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// K base offset for this WG. All 8 K-rows × 4 Q-heads share this K-head.
const ulong offset_src0_base = (i02) * nb02 + (i03 / r3) * nb03;
#pragma unroll
for (int dr = 0; dr < N_K_ROWS_GQA; ++dr) {
const int r0 = r0_base + dr;
const ulong offset_src0 = r0 * nb01 + offset_src0_base;
global half4 * x4 = (global half4 *)(src0 + offset_src0);
float sumf = 0.0f;
#pragma unroll
for (int t = 0; t < 4; ++t) {
const int i = lane_q + t * LANES_PER_QH; // 8, 16, 24-step
const half4 k4 = x4[i];
const float4 q = q_loc[q_id * DK_VEC_GQA + i];
sumf += convert_float(k4.s0) * q.s0
+ convert_float(k4.s1) * q.s1
+ convert_float(k4.s2) * q.s2
+ convert_float(k4.s3) * q.s3;
}
sumf += sub_group_shuffle_xor(sumf, 4);
sumf += sub_group_shuffle_xor(sumf, 2);
sumf += sub_group_shuffle_xor(sumf, 1);
if (lane_q == 0) {
const int im_out = i03 * ne12 + (q_head_lo + q_id);
dst[im_out * ne1 * ne0 + r0] = sumf;
}
}
}
#define N_DV_ROWS_Y8GQA 8
#define GQA_RATIO_Y8GQA 8
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_y8_gqa(
global char * src0,
ulong offset0,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb00,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src0 = (global char *)((global char *)src0 + offset0);
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int sgs_sz = get_max_sub_group_size();
const int r0_base = get_group_id(0) * N_DV_ROWS_Y8GQA;
const int im_kv = get_group_id(2);
const int i02 = im_kv % ne02; // K-head index
const int i03 = im_kv / ne02; // n13 batch index
// GQA Q-heads sharing this K-head.
const int q_head_lo = i02 * GQA_RATIO_Y8GQA;
global float4 * y4_q[GQA_RATIO_Y8GQA];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
const int qh_idx = q_head_lo + qh;
y4_q[qh] = (global float4 *)(src1 + qh_idx * nb12 + i03 * nb13);
}
global half4 * x4_o[N_DV_ROWS_Y8GQA];
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
const int r0 = r0_base + o;
const int r0c = (r0 < ne01) ? r0 : 0;
const ulong off = r0c * nb01 + (i02) * nb02 + (i03 / r3) * nb03;
x4_o[o] = (global half4 *)(src0 + off);
}
float sum[N_DV_ROWS_Y8GQA][GQA_RATIO_Y8GQA] = { {0.0f} };
for (int i = sgs_lid; i < ne00 / 4; i += sgs_sz) {
// load 8 V values (one per DV row), same K-head, K-pos = i.
half4 v[N_DV_ROWS_Y8GQA];
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
v[o] = x4_o[o][i];
}
// load 8 softmax values (one per Q-head).
float4 q[GQA_RATIO_Y8GQA];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
q[qh] = y4_q[qh][i];
}
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
const float4 vf = (float4)(convert_float(v[o].s0),
convert_float(v[o].s1),
convert_float(v[o].s2),
convert_float(v[o].s3));
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
sum[o][qh] += vf.s0 * q[qh].s0
+ vf.s1 * q[qh].s1
+ vf.s2 * q[qh].s2
+ vf.s3 * q[qh].s3;
}
}
}
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
const int r0 = r0_base + o;
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
const float s = sub_group_reduce_add(sum[o][qh]);
if (sgs_lid == 0 && r0 < ne01) {
const int im_out = i03 * ne12 + (q_head_lo + qh);
dst[im_out * ne1 * ne0 + r0] = s;
}
}
}
}
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_x8_gqa4_img(
__read_only image1d_buffer_t src0_img,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int q_id = sgs_lid >> 3; // 0..7: which Q-head (8 per WG)
const int lane_q = sgs_lid & 7; // 0..7: lane within Q-head partition
const int r0_base = get_group_id(0) * N_K_ROWS_GQA;
const int im_kv = get_group_id(2);
const int i02 = im_kv % ne02;
const int i03 = im_kv / ne02;
const int q_head_lo = i02 * GQA_RATIO_GQA;
__local float4 q_loc[GQA_RATIO_GQA * DK_VEC_GQA];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_GQA; ++qh) {
const int qh_idx = q_head_lo + qh;
global float4 * y4 = (global float4 *)(src1 + qh_idx * nb12 + i03 * nb13);
if (sgs_lid < DK_VEC_GQA) {
q_loc[qh * DK_VEC_GQA + sgs_lid] = y4[sgs_lid];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
const int pitch_px_row = (int)(nb01 >> 4);
const int pitch_px_head = (int)(nb02 >> 4);
const int pitch_px_n13 = (int)(nb03 >> 4);
const int head_px_base = i02 * pitch_px_head + (i03 / r3) * pitch_px_n13;
#pragma unroll
for (int dr = 0; dr < N_K_ROWS_GQA; ++dr) {
const int r0 = r0_base + dr;
const int row_px_base = r0 * pitch_px_row + head_px_base;
float sumf = 0.0f;
#pragma unroll
for (int t = 0; t < 2; ++t) {
const int p = lane_q + t * LANES_PER_QH; // pixel idx in row, 0..15
const half8 k8 = as_half8(read_imagef(src0_img, row_px_base + p));
const int i0 = 2 * p; // first half4 idx
const float4 qa = q_loc[q_id * DK_VEC_GQA + i0 ];
const float4 qb = q_loc[q_id * DK_VEC_GQA + i0 + 1];
sumf += convert_float(k8.s0) * qa.s0
+ convert_float(k8.s1) * qa.s1
+ convert_float(k8.s2) * qa.s2
+ convert_float(k8.s3) * qa.s3
+ convert_float(k8.s4) * qb.s0
+ convert_float(k8.s5) * qb.s1
+ convert_float(k8.s6) * qb.s2
+ convert_float(k8.s7) * qb.s3;
}
sumf += sub_group_shuffle_xor(sumf, 4);
sumf += sub_group_shuffle_xor(sumf, 2);
sumf += sub_group_shuffle_xor(sumf, 1);
if (lane_q == 0) {
const int im_out = i03 * ne12 + (q_head_lo + q_id);
dst[im_out * ne1 * ne0 + r0] = sumf;
}
}
}
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_y8_gqa_img(
__read_only image1d_buffer_t src0_img,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int sgs_sz = get_max_sub_group_size();
const int r0_base = get_group_id(0) * N_DV_ROWS_Y8GQA;
const int im_kv = get_group_id(2);
const int i02 = im_kv % ne02;
const int i03 = im_kv / ne02;
const int q_head_lo = i02 * GQA_RATIO_Y8GQA;
// Q (= softmax(KQ)) base pointers per Q-head
global float4 * y4_q[GQA_RATIO_Y8GQA];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
const int qh_idx = q_head_lo + qh;
y4_q[qh] = (global float4 *)(src1 + qh_idx * nb12 + i03 * nb13);
}
const int pitch_px_row = (int)(nb01 >> 3);
const int pitch_px_head = (int)(nb02 >> 3);
const int pitch_px_n13 = (int)(nb03 >> 3);
const int head_px_base = i02 * pitch_px_head + (i03 / r3) * pitch_px_n13;
// per-DV-row pixel base
int row_px_base[N_DV_ROWS_Y8GQA];
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
const int r0 = r0_base + o;
const int r0c = (r0 < ne01) ? r0 : 0;
row_px_base[o] = r0c * pitch_px_row + head_px_base;
}
float sum[N_DV_ROWS_Y8GQA][GQA_RATIO_Y8GQA] = { {0.0f} };
for (int i = sgs_lid; i < ne00 / 4; i += sgs_sz) {
half4 v[N_DV_ROWS_Y8GQA];
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
v[o] = read_imageh(src0_img, row_px_base[o] + i);
}
float4 q[GQA_RATIO_Y8GQA];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
q[qh] = y4_q[qh][i];
}
// 64 mads.
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
const float4 vf = (float4)(convert_float(v[o].s0),
convert_float(v[o].s1),
convert_float(v[o].s2),
convert_float(v[o].s3));
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
sum[o][qh] += vf.s0 * q[qh].s0
+ vf.s1 * q[qh].s1
+ vf.s2 * q[qh].s2
+ vf.s3 * q[qh].s3;
}
}
}
#pragma unroll
for (int o = 0; o < N_DV_ROWS_Y8GQA; ++o) {
const int r0 = r0_base + o;
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_Y8GQA; ++qh) {
const float s = sub_group_reduce_add(sum[o][qh]);
if (sgs_lid == 0 && r0 < ne01) {
const int im_out = i03 * ne12 + (q_head_lo + qh);
dst[im_out * ne1 * ne0 + r0] = s;
}
}
}
}
#define N_K_ROWS_GQA_R4 16
#define GQA_RATIO_R4 4
#define LANES_PER_QH_R4 16 // = 64 / GQA_RATIO_R4
#define DK_VEC_R4 32 // DK / 4 for DK=128
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_x8_gqa_r4_img(
__read_only image1d_buffer_t src0_img,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int q_id = sgs_lid >> 4; // 0..3
const int lane_q = sgs_lid & 15; // 0..15
const int r0_base = get_group_id(0) * N_K_ROWS_GQA_R4;
const int im_kv = get_group_id(2);
const int i02 = im_kv % ne02;
const int i03 = im_kv / ne02;
const int q_head_lo = i02 * GQA_RATIO_R4;
__local float4 q_loc[GQA_RATIO_R4 * DK_VEC_R4];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_R4; ++qh) {
const int qh_idx = q_head_lo + qh;
global float4 * y4 = (global float4 *)(src1 + qh_idx * nb12 + i03 * nb13);
if (sgs_lid < DK_VEC_R4) {
q_loc[qh * DK_VEC_R4 + sgs_lid] = y4[sgs_lid];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
const int pitch_px_row = (int)(nb01 >> 4);
const int pitch_px_head = (int)(nb02 >> 4);
const int pitch_px_n13 = (int)(nb03 >> 4);
const int head_px_base = i02 * pitch_px_head + (i03 / r3) * pitch_px_n13;
#pragma unroll
for (int dr = 0; dr < N_K_ROWS_GQA_R4; ++dr) {
const int r0 = r0_base + dr;
const int row_px_base = r0 * pitch_px_row + head_px_base;
const int p = lane_q;
const half8 k8 = as_half8(read_imagef(src0_img, row_px_base + p));
const int i0 = 2 * p;
const float4 qa = q_loc[q_id * DK_VEC_R4 + i0 ];
const float4 qb = q_loc[q_id * DK_VEC_R4 + i0 + 1];
float sumf =
convert_float(k8.s0) * qa.s0
+ convert_float(k8.s1) * qa.s1
+ convert_float(k8.s2) * qa.s2
+ convert_float(k8.s3) * qa.s3
+ convert_float(k8.s4) * qb.s0
+ convert_float(k8.s5) * qb.s1
+ convert_float(k8.s6) * qb.s2
+ convert_float(k8.s7) * qb.s3;
sumf += sub_group_shuffle_xor(sumf, 8);
sumf += sub_group_shuffle_xor(sumf, 4);
sumf += sub_group_shuffle_xor(sumf, 2);
sumf += sub_group_shuffle_xor(sumf, 1);
if (lane_q == 0) {
const int im_out = i03 * ne12 + (q_head_lo + q_id);
dst[im_out * ne1 * ne0 + r0] = sumf;
}
}
}
#define N_K_ROWS_GQA_R2_DK256 16
#define GQA_RATIO_R2 2
#define LANES_PER_QH_R2 32 // = 64 / GQA_RATIO_R2
#define DK_VEC_DK256 64 // DK / 4 for DK=256
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mat_f16_f32_l4_x8_gqa_r2_dk256_img(
__read_only image1d_buffer_t src0_img,
global char * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
ulong nb01,
ulong nb02,
ulong nb03,
int ne10,
int ne11,
int ne12,
ulong nb10,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src1 = (global char *)((global char *)src1 + offset1);
dst = (global float*)((global char *)dst + offsetd);
const int sgs_lid = get_sub_group_local_id();
const int q_id = sgs_lid >> 5; // 0..1
const int lane_q = sgs_lid & 31; // 0..31
const int r0_base = get_group_id(0) * N_K_ROWS_GQA_R2_DK256;
const int im_kv = get_group_id(2);
const int i02 = im_kv % ne02;
const int i03 = im_kv / ne02;
const int q_head_lo = i02 * GQA_RATIO_R2;
__local float4 q_loc[GQA_RATIO_R2 * DK_VEC_DK256];
#pragma unroll
for (int qh = 0; qh < GQA_RATIO_R2; ++qh) {
const int qh_idx = q_head_lo + qh;
global float4 * y4 = (global float4 *)(src1 + qh_idx * nb12 + i03 * nb13);
q_loc[qh * DK_VEC_DK256 + sgs_lid] = y4[sgs_lid];
}
barrier(CLK_LOCAL_MEM_FENCE);
const int pitch_px_row = (int)(nb01 >> 4);
const int pitch_px_head = (int)(nb02 >> 4);
const int pitch_px_n13 = (int)(nb03 >> 4);
const int head_px_base = i02 * pitch_px_head + (i03 / r3) * pitch_px_n13;
#pragma unroll
for (int dr = 0; dr < N_K_ROWS_GQA_R2_DK256; ++dr) {
const int r0 = r0_base + dr;
const int row_px_base = r0 * pitch_px_row + head_px_base;
const int p = lane_q;
const half8 k8 = as_half8(read_imagef(src0_img, row_px_base + p));
const int i0 = 2 * p;
const float4 qa = q_loc[q_id * DK_VEC_DK256 + i0 ];
const float4 qb = q_loc[q_id * DK_VEC_DK256 + i0 + 1];
float sumf =
convert_float(k8.s0) * qa.s0
+ convert_float(k8.s1) * qa.s1
+ convert_float(k8.s2) * qa.s2
+ convert_float(k8.s3) * qa.s3
+ convert_float(k8.s4) * qb.s0
+ convert_float(k8.s5) * qb.s1
+ convert_float(k8.s6) * qb.s2
+ convert_float(k8.s7) * qb.s3;
sumf += sub_group_shuffle_xor(sumf, 16);
sumf += sub_group_shuffle_xor(sumf, 8);
sumf += sub_group_shuffle_xor(sumf, 4);
sumf += sub_group_shuffle_xor(sumf, 2);
sumf += sub_group_shuffle_xor(sumf, 1);
if (lane_q == 0) {
const int im_out = i03 * ne12 + (q_head_lo + q_id);
dst[im_out * ne1 * ne0 + r0] = sumf;
}
}
}
+1
View File
@@ -14,6 +14,7 @@
#define GGML_SYCL_BACKEND_HPP
#include "binbcast.hpp"
#include "col2im-1d.hpp"
#include "common.hpp"
#include "concat.hpp"
#include "conv.hpp"
+102
View File
@@ -0,0 +1,102 @@
#include "col2im-1d.hpp"
template <typename T>
static void col2im_1d_sycl(
const T * col,
T * dst,
const int T_in,
const sycl::uint3 T_out_fd,
const int K,
const int K_OC,
const int32_t s0,
const int32_t p0,
const int total,
dpct::queue_ptr stream) {
const uint32_t block_size = SYCL_COL2IM_1D_BLOCK_SIZE;
const uint32_t num_blocks = (uint32_t) ((total + block_size - 1) / block_size);
stream->parallel_for(
sycl::nd_range<3>(
sycl::range<3>(1, 1, num_blocks * block_size),
sycl::range<3>(1, 1, block_size)),
[=](sycl::nd_item<3> item_ct1) {
const int idx = (int) item_ct1.get_global_id(2);
if (idx >= total) {
return;
}
const sycl::uint2 qr = fast_div_modulo((uint32_t) idx, T_out_fd);
const int oc = (int) qr.x();
const int t_out = (int) qr.y();
const int t_abs = t_out + p0;
int t_in_min = (t_abs - K + s0) / s0;
if (t_in_min < 0) {
t_in_min = 0;
}
int t_in_max = t_abs / s0;
if (t_in_max >= T_in) {
t_in_max = T_in - 1;
}
float sum = 0.0f;
for (int t_in = t_in_min; t_in <= t_in_max; ++t_in) {
const int k = t_abs - t_in * s0;
sum += static_cast<float>(col[(oc * K + k) + t_in * K_OC]);
}
dst[idx] = static_cast<T>(sum);
});
}
void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
GGML_ASSERT(src0 != nullptr);
GGML_ASSERT(ggml_is_contiguous(src0));
GGML_ASSERT(src0->type == dst->type);
const int32_t s0 = ((const int32_t *) dst->op_params)[0];
const int32_t OC = ((const int32_t *) dst->op_params)[1];
const int32_t p0 = ((const int32_t *) dst->op_params)[2];
const int K_OC = (int) src0->ne[0];
const int T_in = (int) src0->ne[1];
const int K = K_OC / OC;
const int T_out = (int) dst->ne[0];
GGML_ASSERT(OC > 0);
GGML_ASSERT(K_OC % OC == 0);
const sycl::uint3 T_out_fd = init_fastdiv_values((uint32_t) T_out);
const int total = T_out * OC;
dpct::queue_ptr stream = ctx.stream();
switch (src0->type) {
case GGML_TYPE_F32:
col2im_1d_sycl<float>(
(const float *) src0->data,
(float *) dst->data,
T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
break;
case GGML_TYPE_F16:
col2im_1d_sycl<sycl::half>(
(const sycl::half *) src0->data,
(sycl::half *) dst->data,
T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
break;
#ifdef GGML_SYCL_HAS_BF16
case GGML_TYPE_BF16:
col2im_1d_sycl<sycl::ext::oneapi::bfloat16>(
(const sycl::ext::oneapi::bfloat16 *) src0->data,
(sycl::ext::oneapi::bfloat16 *) dst->data,
T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
break;
#endif
default:
GGML_ABORT("col2im_1d: unsupported type %d", src0->type);
}
}
+8
View File
@@ -0,0 +1,8 @@
#ifndef GGML_SYCL_COL2IM_1D_HPP
#define GGML_SYCL_COL2IM_1D_HPP
#include "common.hpp"
void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
#endif // GGML_SYCL_COL2IM_1D_HPP
+1 -1
View File
@@ -59,7 +59,7 @@ void ggml_sycl_host_free(void* ptr);
extern int g_ggml_sycl_debug;
extern int g_ggml_sycl_disable_optimize;
extern int g_ggml_sycl_enable_optimize;
extern int g_ggml_sycl_prioritize_dmmv;
extern int g_ggml_sycl_enable_flash_attention;
extern int g_ggml_sycl_dev2dev_memcpy;
+706
View File
@@ -1,6 +1,7 @@
#include "cpy.hpp"
#include <float.h>
#include <vector>
#include "dequantize.hpp"
#include "ggml-sycl/common.hpp"
@@ -50,6 +51,57 @@ static void cpy_1_i32_i32(const char * cxi, char * cdsti) {
*dsti = *xi;
}
static void cpy_1_f32_i32(const char * cxi, char * cdsti) {
const float * xi = (const float *) cxi;
int32_t * dsti = (int32_t *) cdsti;
*dsti = (int32_t) *xi;
}
static void cpy_1_i32_f32(const char * cxi, char * cdsti) {
const int32_t * xi = (const int32_t *) cxi;
float * dsti = (float *) cdsti;
*dsti = (float) *xi;
}
#ifdef GGML_SYCL_HAS_BF16
static void cpy_1_f32_bf16(const char * cxi, char * cdsti) {
const float * xi = (const float *) cxi;
sycl::ext::oneapi::bfloat16 * dsti = (sycl::ext::oneapi::bfloat16 *) cdsti;
*dsti = sycl::ext::oneapi::bfloat16(*xi);
}
static void cpy_1_bf16_f32(const char * cxi, char * cdsti) {
const sycl::ext::oneapi::bfloat16 * xi = (const sycl::ext::oneapi::bfloat16 *) cxi;
float * dsti = (float *) cdsti;
*dsti = static_cast<float>(*xi);
}
static void cpy_1_bf16_bf16(const char * cxi, char * cdsti) {
const sycl::ext::oneapi::bfloat16 * xi = (const sycl::ext::oneapi::bfloat16 *) cxi;
sycl::ext::oneapi::bfloat16 * dsti = (sycl::ext::oneapi::bfloat16 *) cdsti;
*dsti = *xi;
}
static void cpy_1_f16_bf16(const char * cxi, char * cdsti) {
const sycl::half * xi = (const sycl::half *) cxi;
sycl::ext::oneapi::bfloat16 * dsti = (sycl::ext::oneapi::bfloat16 *) cdsti;
*dsti = sycl::ext::oneapi::bfloat16(static_cast<float>(*xi));
}
static void cpy_1_bf16_f16(const char * cxi, char * cdsti) {
const sycl::ext::oneapi::bfloat16 * xi = (const sycl::ext::oneapi::bfloat16 *) cxi;
sycl::half * dsti = (sycl::half *) cdsti;
*dsti = sycl::half(static_cast<float>(*xi));
}
#endif
template <cpy_kernel_t cpy_1>
static void cpy_f32_f16(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, const int ne02,
const int nb00, const int nb01, const int nb02, const int nb03, const int ne10, const int ne11,
@@ -247,6 +299,38 @@ static void ggml_cpy_f32_f16_sycl(const char * cx, char * cdst, const int ne, co
}
}
static void ggml_cpy_f32_i32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
{
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_f32_i32>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
}
static void ggml_cpy_i32_f32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
{
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_i32_f32>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
}
static void ggml_cpy_f32_q8_0_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
@@ -376,6 +460,19 @@ static void ggml_cpy_q5_1_f32_sycl(const char * cx, char * cdst, const int ne, c
});
}
static void ggml_cpy_mxfp4_f32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ne;
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_f32<cpy_blck_q_f32<dequantize_mxfp4, QK_MXFP4>, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_f32_iq4_nl_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
@@ -389,6 +486,269 @@ static void ggml_cpy_f32_iq4_nl_sycl(const char * cx, char * cdst, const int ne,
});
}
static void cpy_blck_f16_q4_0(const char * cxi, char * cdsti) {
const sycl::half * xi = (const sycl::half *) cxi;
float xf[QK4_0];
for (int j = 0; j < QK4_0; ++j) {
xf[j] = (float) xi[j];
}
cpy_blck_f32_q4_0((const char *) xf, cdsti);
}
static void cpy_blck_f16_q4_1(const char * cxi, char * cdsti) {
const sycl::half * xi = (const sycl::half *) cxi;
float xf[QK4_1];
for (int j = 0; j < QK4_1; ++j) {
xf[j] = (float) xi[j];
}
cpy_blck_f32_q4_1((const char *) xf, cdsti);
}
static void cpy_blck_f16_q5_0(const char * cxi, char * cdsti) {
const sycl::half * xi = (const sycl::half *) cxi;
float xf[QK5_0];
for (int j = 0; j < QK5_0; ++j) {
xf[j] = (float) xi[j];
}
cpy_blck_f32_q5_0((const char *) xf, cdsti);
}
static void ggml_cpy_f16_q4_0_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
GGML_ASSERT(ne % QK4_0 == 0);
const int num_blocks = ne / QK4_0;
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_q<cpy_blck_f16_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_f16_q4_1_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
GGML_ASSERT(ne % QK4_1 == 0);
const int num_blocks = ne / QK4_1;
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_q<cpy_blck_f16_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_f16_q5_0_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
GGML_ASSERT(ne % QK5_0 == 0);
const int num_blocks = ne / QK5_0;
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_q<cpy_blck_f16_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03,
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static bool ggml_sycl_is_quantized_type(enum ggml_type type) {
switch (type) {
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_MXFP4:
case GGML_TYPE_NVFP4:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
case GGML_TYPE_IQ2_XXS:
case GGML_TYPE_IQ2_XS:
case GGML_TYPE_IQ2_S:
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ3_S:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
return true;
default:
return false;
}
}
static bool ggml_sycl_can_quantize_rows_sycl(enum ggml_type type) {
switch (type) {
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_MXFP4:
case GGML_TYPE_NVFP4:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
return true;
default:
return false;
}
}
template <typename SrcScalar>
static inline float ggml_sycl_src_to_f32(const SrcScalar & x) {
return (float) x;
}
#ifdef GGML_SYCL_HAS_BF16
template <>
inline float ggml_sycl_src_to_f32<sycl::ext::oneapi::bfloat16>(const sycl::ext::oneapi::bfloat16 & x) {
return static_cast<float>(x);
}
template <>
inline float ggml_sycl_src_to_f32<ggml_bf16_t>(const ggml_bf16_t & x) {
union {
uint32_t u32;
float f32;
} value;
value.u32 = (uint32_t) x.bits << 16;
return value.f32;
}
#endif
template <typename SrcScalar, cpy_kernel_t quantize_block, int qk>
static void ggml_sycl_quantize_rows_q(const char * cx, char * cdst, const int64_t ne,
const int64_t ne00, const int64_t ne01, const int64_t ne02,
const size_t nb00, const size_t nb01, const size_t nb02, const size_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12,
const size_t nb10, const size_t nb11, const size_t nb12, const size_t nb13,
queue_ptr stream) {
GGML_ASSERT(ne % qk == 0);
GGML_ASSERT(ne00 % qk == 0);
const int64_t total_blocks = ne / qk;
constexpr int block_size = 256;
const int64_t grid_size = ceil_div(total_blocks, (int64_t) block_size);
stream->parallel_for(sycl::nd_range<1>(grid_size * block_size, block_size), [=](sycl::nd_item<1> item_ct1) {
const int64_t block_idx = item_ct1.get_global_linear_id();
if (block_idx >= total_blocks) {
return;
}
const int64_t i = block_idx * qk;
const int64_t i03 = i / (ne00 * ne01 * ne02);
const int64_t i02 = (i - i03 * ne00 * ne01 * ne02) / (ne00 * ne01);
const int64_t i01 = (i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00) / ne00;
const int64_t i00 = i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00 - i01 * ne00;
const size_t x_offset = i00 * nb00 + i01 * nb01 + i02 * nb02 + i03 * nb03;
const int64_t i13 = i / (ne10 * ne11 * ne12);
const int64_t i12 = (i - i13 * ne10 * ne11 * ne12) / (ne10 * ne11);
const int64_t i11 = (i - i13 * ne10 * ne11 * ne12 - i12 * ne10 * ne11) / ne10;
const int64_t i10 = i - i13 * ne10 * ne11 * ne12 - i12 * ne10 * ne11 - i11 * ne10;
const size_t dst_offset = (i10 / qk) * nb10 + i11 * nb11 + i12 * nb12 + i13 * nb13;
float xf[qk];
if (nb00 == sizeof(SrcScalar)) {
const SrcScalar * src_row = (const SrcScalar *) (cx + x_offset);
for (int j = 0; j < qk; ++j) {
xf[j] = ggml_sycl_src_to_f32(src_row[j]);
}
} else {
for (int j = 0; j < qk; ++j) {
const SrcScalar * src_val = (const SrcScalar *) (cx + x_offset + j * nb00);
xf[j] = ggml_sycl_src_to_f32(*src_val);
}
}
quantize_block((const char *) xf, cdst + dst_offset);
});
}
template <typename SrcScalar>
static void ggml_sycl_quantize_rows_sycl(const char * cx, char * cdst, const ggml_tensor * src0, const ggml_tensor * src1,
const int64_t ne, const int64_t ne00, const int64_t ne01, const int64_t ne02,
const size_t nb00, const size_t nb01, const size_t nb02, const size_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const size_t nb10,
const size_t nb11, const size_t nb12, const size_t nb13, queue_ptr stream) {
GGML_UNUSED(src0);
GGML_UNUSED(src1);
switch (src1->type) {
case GGML_TYPE_Q8_0:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q8_0, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, stream);
break;
case GGML_TYPE_Q1_0:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q1_0, QK1_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, stream);
break;
case GGML_TYPE_Q5_1:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q5_1, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, stream);
break;
case GGML_TYPE_Q5_0:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, stream);
break;
case GGML_TYPE_Q4_1:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, stream);
break;
case GGML_TYPE_Q4_0:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, stream);
break;
case GGML_TYPE_IQ4_NL:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_iq4_nl, QK4_NL>(cx, cdst, ne, ne00, ne01, ne02, nb00,
nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, stream);
break;
case GGML_TYPE_MXFP4:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_mxfp4, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, stream);
break;
case GGML_TYPE_NVFP4:
ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_nvfp4, QK_NVFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, stream);
break;
default:
GGML_ABORT("unsupported quantized target type in sycl quantizer src1->type=%s\n",
ggml_type_name(src1->type));
}
}
static void ggml_cpy_f16_f16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
@@ -509,8 +869,269 @@ static void ggml_cpy_q4_1_q4_1(const char * cx, char * cdst, const int ne, const
});
}
static void ggml_cpy_q1_0_q1_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_q1_0, QK1_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_mxfp4_mxfp4(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_mxfp4, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_nvfp4_nvfp4(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_nvfp4, QK_NVFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_q2_K_q2_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_q2_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_q3_K_q3_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_q3_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_q4_K_q4_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_q4_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_q5_K_q5_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_q5_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_q6_K_q6_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_q6_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq2_xxs_iq2_xxs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq2_xxs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq2_xs_iq2_xs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq2_xs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq2_s_iq2_s(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq2_s, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq3_xxs_iq3_xxs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq3_xxs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq1_s_iq1_s(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq1_s, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq1_m_iq1_m(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq1_m, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq4_nl_iq4_nl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq4_nl, QK4_NL>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq3_s_iq3_s(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq3_s, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_iq4_xs_iq4_xs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
cpy_q_q<block_iq4_xs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
});
}
#ifdef GGML_SYCL_HAS_BF16
static void ggml_cpy_f32_bf16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_f32_bf16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_bf16_f32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_bf16_f32>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_bf16_bf16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_bf16_bf16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_f16_bf16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_f16_bf16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
static void ggml_cpy_bf16_f16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
const int nb12, const int nb13, queue_ptr stream) {
const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
[=](sycl::nd_item<3> item_ct1) {
cpy_f32_f16<cpy_1_bf16_f16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, item_ct1);
});
}
#endif
void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1) try {
// Unlike other operators ggml_sycl_cpy takes 2 distinct tensors instead of a dst ggml_tensor and rely on its src field
GGML_SYCL_DEBUG("ggml_sycl_cpy: src0->type=%s, src1->type=%s\n",
ggml_type_name(src0->type), ggml_type_name(src1->type));
scope_op_debug_print scope_dbg_print(__func__, src1, /*num_src=*/0, debug_get_tensor_str("\tsrc0", src0));
const int64_t ne = ggml_nelements(src0);
GGML_ASSERT(ne == ggml_nelements(src1));
@@ -525,12 +1146,31 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
if ((src0->type == src1->type) && (ggml_is_contiguous(src0) && ggml_is_contiguous(src1))) {
GGML_SYCL_DEBUG("%s: memcpy path\n", __func__);
main_stream->memcpy(src1_ddc, src0_ddc, ggml_nbytes(src0));
} else if (src0->type == GGML_TYPE_F32 && ggml_sycl_is_quantized_type(src1->type)) {
GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type));
ggml_sycl_quantize_rows_sycl<float>(src0_ddc, src1_ddc, src0, src1, ne, ne00, ne01, ne02, nb00, nb01,
nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F16 && ggml_sycl_is_quantized_type(src1->type)) {
GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type));
ggml_sycl_quantize_rows_sycl<sycl::half>(src0_ddc, src1_ddc, src0, src1, ne, ne00, ne01, ne02, nb00,
nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
main_stream);
#ifdef GGML_SYCL_HAS_BF16
} else if (src0->type == GGML_TYPE_BF16 && ggml_sycl_is_quantized_type(src1->type)) {
GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type));
ggml_sycl_quantize_rows_sycl<ggml_bf16_t>(src0_ddc, src1_ddc, src0, src1, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11,
nb12, nb13, main_stream);
#endif
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
ggml_cpy_f32_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F16) {
ggml_cpy_f32_f16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_I32) {
ggml_cpy_f32_i32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_Q8_0) {
ggml_cpy_f32_q8_0_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
@@ -546,12 +1186,24 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16) {
ggml_cpy_f16_f16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_Q4_0) {
ggml_cpy_f16_q4_0_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_Q4_1) {
ggml_cpy_f16_q4_1_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_Q5_0) {
ggml_cpy_f16_q5_0_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02,
nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_I16 && src1->type == GGML_TYPE_I16) {
ggml_cpy_i16_i16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_I32 && src1->type == GGML_TYPE_I32) {
ggml_cpy_i32_i32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_I32 && src1->type == GGML_TYPE_F32) {
ggml_cpy_i32_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q4_0 && src1->type == GGML_TYPE_F32) {
ggml_cpy_q4_0_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
@@ -573,6 +1225,9 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
} else if (src0->type == GGML_TYPE_Q5_1 && src1->type == GGML_TYPE_F32) {
ggml_cpy_q5_1_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_MXFP4 && src1->type == GGML_TYPE_F32) {
ggml_cpy_mxfp4_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_IQ4_NL) {
ggml_cpy_f32_iq4_nl_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
nb10, nb11, nb12, nb13, main_stream);
@@ -586,6 +1241,57 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
ggml_cpy_q4_0_q4_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q4_1 && src1->type == GGML_TYPE_Q4_1) {
ggml_cpy_q4_1_q4_1(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q1_0 && src1->type == GGML_TYPE_Q1_0) {
ggml_cpy_q1_0_q1_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_MXFP4 && src1->type == GGML_TYPE_MXFP4) {
ggml_cpy_mxfp4_mxfp4(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_NVFP4 && src1->type == GGML_TYPE_NVFP4) {
ggml_cpy_nvfp4_nvfp4(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q2_K && src1->type == GGML_TYPE_Q2_K) {
ggml_cpy_q2_K_q2_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q3_K && src1->type == GGML_TYPE_Q3_K) {
ggml_cpy_q3_K_q3_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q4_K && src1->type == GGML_TYPE_Q4_K) {
ggml_cpy_q4_K_q4_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q5_K && src1->type == GGML_TYPE_Q5_K) {
ggml_cpy_q5_K_q5_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_Q6_K && src1->type == GGML_TYPE_Q6_K) {
ggml_cpy_q6_K_q6_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ2_XXS && src1->type == GGML_TYPE_IQ2_XXS) {
ggml_cpy_iq2_xxs_iq2_xxs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ2_XS && src1->type == GGML_TYPE_IQ2_XS) {
ggml_cpy_iq2_xs_iq2_xs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ2_S && src1->type == GGML_TYPE_IQ2_S) {
ggml_cpy_iq2_s_iq2_s(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ3_XXS && src1->type == GGML_TYPE_IQ3_XXS) {
ggml_cpy_iq3_xxs_iq3_xxs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ1_S && src1->type == GGML_TYPE_IQ1_S) {
ggml_cpy_iq1_s_iq1_s(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ1_M && src1->type == GGML_TYPE_IQ1_M) {
ggml_cpy_iq1_m_iq1_m(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ4_NL && src1->type == GGML_TYPE_IQ4_NL) {
ggml_cpy_iq4_nl_iq4_nl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ3_S && src1->type == GGML_TYPE_IQ3_S) {
ggml_cpy_iq3_s_iq3_s(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_IQ4_XS && src1->type == GGML_TYPE_IQ4_XS) {
ggml_cpy_iq4_xs_iq4_xs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
#ifdef GGML_SYCL_HAS_BF16
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_BF16) {
ggml_cpy_f32_bf16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F32) {
ggml_cpy_bf16_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_BF16) {
ggml_cpy_bf16_bf16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_BF16) {
ggml_cpy_f16_bf16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F16) {
ggml_cpy_bf16_f16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
nb11, nb12, nb13, main_stream);
#endif
} else {
GGML_LOG_ERROR("%s: unsupported type combination (%s to %s)\n", __func__, ggml_type_name(src0->type),
ggml_type_name(src1->type));
+1 -1
View File
@@ -317,7 +317,7 @@ inline void cpy_blck_f32_nvfp4(const char * cxi, char * cdsti) {
const uint8_t ue = ggml_fp32_to_ue4m3(amax / 6.0f);
dsti->d[s] = ue;
const float d = ggml_ue4m3_to_fp32(ue);
const float d = ggml_sycl_ue4m3_to_fp32(ue);
for (int j = 0; j < QK_NVFP4_SUB / 2; ++j) {
const uint8_t x0 = best_index_mxfp4(xb[0 + j], d);
+255
View File
@@ -0,0 +1,255 @@
#include "cross_entropy_loss.hpp"
#include <cstdint>
#include <cmath>
template <bool has_shared>
static __dpct_inline__ void cross_entropy_loss_f32_kernel(
const float * __restrict__ logits,
const float * __restrict__ labels,
float * __restrict__ row_loss,
const int nclasses,
const int nrows,
float * __restrict__ smem,
const sycl::nd_item<3> & item) {
const int row = item.get_group(2);
const int tid = item.get_local_id(2);
logits += (int64_t) row * nclasses;
labels += (int64_t) row * nclasses;
float max_logit = -INFINITY;
for (int i = tid; i < nclasses; i += WARP_SIZE) {
const float v = logits[i];
max_logit = sycl::fmax(max_logit, v);
if (has_shared) {
smem[i] = v;
}
}
max_logit = warp_reduce_max<WARP_SIZE>(max_logit);
float sum_exp = 0.0f;
for (int i = tid; i < nclasses; i += WARP_SIZE) {
const float v = has_shared ? smem[i] : logits[i];
sum_exp += sycl::exp(v - max_logit);
}
sum_exp = warp_reduce_sum<WARP_SIZE>(sum_exp);
const float log_sum = sycl::log(sum_exp);
float loss = 0.0f;
for (int i = tid; i < nclasses; i += WARP_SIZE) {
const float v = has_shared ? smem[i] : logits[i];
loss += (v - max_logit - log_sum) * labels[i];
}
loss = -warp_reduce_sum<WARP_SIZE>(loss) / (float) nrows;
if (tid == 0) {
row_loss[row] = loss;
}
}
template <bool has_shared>
static __dpct_inline__ void cross_entropy_loss_back_f32_kernel(
const float * __restrict__ grad,
const float * __restrict__ logits,
const float * __restrict__ labels,
float * __restrict__ dst,
const int nclasses,
const int nrows,
float * __restrict__ smem,
const sycl::nd_item<3> & item) {
const int row = item.get_group(2);
const int tid = item.get_local_id(2);
logits += (int64_t) row * nclasses;
labels += (int64_t) row * nclasses;
dst += (int64_t) row * nclasses;
float max_logit = -INFINITY;
for (int i = tid; i < nclasses; i += WARP_SIZE) {
const float v = logits[i];
max_logit = sycl::fmax(max_logit, v);
if (has_shared) {
smem[i] = v;
}
}
max_logit = warp_reduce_max<WARP_SIZE>(max_logit);
float sum_exp = 0.0f;
for (int i = tid; i < nclasses; i += WARP_SIZE) {
const float v = sycl::exp((has_shared ? smem[i] : logits[i]) - max_logit);
sum_exp += v;
if (has_shared) {
smem[i] = v;
} else {
dst[i] = v;
}
}
sum_exp = warp_reduce_sum<WARP_SIZE>(sum_exp);
const float inv_sum = 1.0f / sum_exp;
const float d_by_nrows = grad[0] / (float) nrows;
for (int i = tid; i < nclasses; i += WARP_SIZE) {
const float sm_num = has_shared ? smem[i] : dst[i];
dst[i] = (sm_num * inv_sum - labels[i]) * d_by_nrows;
}
}
static void cross_entropy_reduce_rows(
ggml_backend_sycl_context & ctx,
const float * row_loss,
float * dst,
const int64_t nrows) {
if (nrows == 1) {
SYCL_CHECK(CHECK_TRY_ERROR(
ctx.stream()->memcpy(dst, row_loss, sizeof(float))));
return;
}
ggml_sycl_pool_alloc<float> tmp_alloc(ctx.pool(), nrows);
float * tmp = tmp_alloc.get();
SYCL_CHECK(CHECK_TRY_ERROR(
ctx.stream()->memcpy(tmp, row_loss, nrows * sizeof(float))));
int64_t cur = nrows;
while (cur > 1) {
const int64_t out = (cur + WARP_SIZE - 1) / WARP_SIZE;
const sycl::range<3> block(1, 1, WARP_SIZE);
const sycl::range<3> grid(1, 1, out);
ctx.stream()->parallel_for(
sycl::nd_range<3>(grid * block, block),
[=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
const int row = item.get_group(2);
const int tid = item.get_local_id(2);
const int64_t i = (int64_t) row * WARP_SIZE + tid;
float v = i < cur ? tmp[i] : 0.0f;
v = warp_reduce_sum<WARP_SIZE>(v);
if (tid == 0) {
tmp[row] = v;
}
});
cur = out;
}
SYCL_CHECK(CHECK_TRY_ERROR(
ctx.stream()->memcpy(dst, tmp, sizeof(float))));
}
void ggml_sycl_cross_entropy_loss(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
GGML_ASSERT(src0->type == GGML_TYPE_F32);
GGML_ASSERT(src1->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous(src0));
GGML_ASSERT(ggml_is_contiguous(src1));
GGML_ASSERT(ggml_is_contiguous(dst));
GGML_ASSERT(ggml_are_same_shape(src0, src1));
GGML_ASSERT(ggml_is_scalar(dst));
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
const int64_t nclasses = src0->ne[0];
const int64_t nrows = ggml_nrows(src0);
const float * logits_d = (const float *) src0->data;
const float * labels_d = (const float *) src1->data;
float * dst_d = (float *) dst->data;
ggml_sycl_pool_alloc<float> row_loss_alloc(ctx.pool(), nrows);
float * row_loss = row_loss_alloc.get();
const sycl::range<3> block(1, 1, WARP_SIZE);
const sycl::range<3> grid(1, 1, nrows);
const size_t nbytes_shared = (size_t) nclasses * sizeof(float);
const size_t smpbo = ggml_sycl_info().devices[ctx.device].smpbo;
if (nbytes_shared <= smpbo) {
ctx.stream()->submit([&](sycl::handler & cgh) {
sycl::local_accessor<float, 1> smem(sycl::range<1>(nclasses), cgh);
cgh.parallel_for(
sycl::nd_range<3>(grid * block, block),
[=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
cross_entropy_loss_f32_kernel<true>(
logits_d, labels_d, row_loss,
(int) nclasses, (int) nrows,
get_pointer(smem), item);
});
});
} else {
ctx.stream()->parallel_for(
sycl::nd_range<3>(grid * block, block),
[=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
cross_entropy_loss_f32_kernel<false>(
logits_d, labels_d, row_loss,
(int) nclasses, (int) nrows,
nullptr, item);
});
}
cross_entropy_reduce_rows(ctx, row_loss, dst_d, nrows);
}
void ggml_sycl_cross_entropy_loss_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/3);
const ggml_tensor * grad = dst->src[0];
const ggml_tensor * src0f = dst->src[1];
const ggml_tensor * src1f = dst->src[2];
GGML_ASSERT(grad->type == GGML_TYPE_F32);
GGML_ASSERT(src0f->type == GGML_TYPE_F32);
GGML_ASSERT(src1f->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_scalar(grad));
GGML_ASSERT(ggml_is_contiguous(grad));
GGML_ASSERT(ggml_is_contiguous(src0f));
GGML_ASSERT(ggml_is_contiguous(src1f));
GGML_ASSERT(ggml_is_contiguous(dst));
GGML_ASSERT(ggml_are_same_shape(src0f, src1f));
GGML_ASSERT(ggml_are_same_shape(src0f, dst));
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
const int64_t nclasses = src0f->ne[0];
const int64_t nrows = ggml_nrows(src0f);
const float * grad_d = (const float *) grad->data;
const float * logits_d = (const float *) src0f->data;
const float * labels_d = (const float *) src1f->data;
float * dst_d = (float *) dst->data;
const sycl::range<3> block(1, 1, WARP_SIZE);
const sycl::range<3> grid(1, 1, nrows);
const size_t nbytes_shared = (size_t) nclasses * sizeof(float);
const size_t smpbo = ggml_sycl_info().devices[ctx.device].smpbo;
if (nbytes_shared <= smpbo) {
ctx.stream()->submit([&](sycl::handler & cgh) {
sycl::local_accessor<float, 1> smem(sycl::range<1>(nclasses), cgh);
cgh.parallel_for(
sycl::nd_range<3>(grid * block, block),
[=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
cross_entropy_loss_back_f32_kernel<true>(
grad_d, logits_d, labels_d, dst_d,
(int) nclasses, (int) nrows,
get_pointer(smem), item);
});
});
} else {
ctx.stream()->parallel_for(
sycl::nd_range<3>(grid * block, block),
[=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
cross_entropy_loss_back_f32_kernel<false>(
grad_d, logits_d, labels_d, dst_d,
(int) nclasses, (int) nrows,
nullptr, item);
});
}
}
@@ -0,0 +1,7 @@
#pragma once
#include "common.hpp"
void ggml_sycl_cross_entropy_loss(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_cross_entropy_loss_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+15 -12
View File
@@ -680,14 +680,14 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
q16[2] = q2[0] & 0x0f0f;
q16[3] = q2[0] & 0xf0f0;
float4 s = {0.f, 0.f, 0.f, 0.f};
sycl::float4 s = {0.f, 0.f, 0.f, 0.f};
float smin = 0;
for (int l = 0; l < 2; ++l) {
s.x += y1[l] * q4[l+0]; s.y += y1[l+32] * q4[l+2];
s.z += y2[l] * q4[l+4]; s.w += y2[l+32] * q4[l+6];
s.x() += y1[l] * q4[l+0]; s.y() += y1[l+32] * q4[l+2];
s.z() += y2[l] * q4[l+4]; s.w() += y2[l+32] * q4[l+6];
smin += y1[l] * sc[2] + y1[l+32] * sc[3] + y2[l] * sc[6] + y2[l+32] * sc[7];
}
tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin;
tmp += dall * (s.x() * sc[0] + s.y() * sc[1] * 1.f/16.f + s.z() * sc[4] + s.w() * sc[5] * 1.f/16.f) - dmin * smin;
#endif
}
@@ -835,14 +835,14 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
q16[2] = q2[0] & 0x0f0f;
q16[3] = q2[0] & 0xf0f0;
float4 s = {0.f, 0.f, 0.f, 0.f};
sycl::float4 s = {0.f, 0.f, 0.f, 0.f};
float smin = 0;
for (int l = 0; l < 2; ++l) {
s.x += y1[l] * q4[l+0]; s.y += y1[l+32] * q4[l+2];
s.z += y2[l] * q4[l+4]; s.w += y2[l+32] * q4[l+6];
s.x() += y1[l] * q4[l+0]; s.y() += y1[l+32] * q4[l+2];
s.z() += y2[l] * q4[l+4]; s.w() += y2[l+32] * q4[l+6];
smin += y1[l] * sc[2] + y1[l+32] * sc[3] + y2[l] * sc[6] + y2[l+32] * sc[7];
}
tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin;
tmp += dall * (s.x() * sc[0] + s.y() * sc[1] * 1.f/16.f + s.z() * sc[4] + s.w() * sc[5] * 1.f/16.f) - dmin * smin;
#endif
}
@@ -1126,7 +1126,7 @@ static void dequantize_mul_mat_vec_q5_k_reorder(const void *__restrict__ vx,
// sum up partial sums and write back result
#pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) {
for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
}
@@ -1762,10 +1762,13 @@ static void dequantize_mul_mat_vec_q5_K_sycl_reorder(const void *vx, const float
const int nrows,
dpct::queue_ptr stream) {
GGML_ASSERT(ncols % QK_K == 0);
const sycl::range<3> block_dims(1, 1, QK_WARP_SIZE);
const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q5_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
});
}
+40 -16
View File
@@ -9,9 +9,12 @@
#define SYCL_LOCAL_ID_CALC(ITEM, IDX) \
(ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX))
static void acc_f32(const float * x, const float * y, float * dst, const int64_t ne,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
const int64_t i = SYCL_LOCAL_ID_CALC(item_ct1, 2);
@@ -30,9 +33,18 @@ static void acc_f32(const float * x, const float * y, float * dst, const int64_t
tmp -= i11 * s11;
const int64_t i10 = tmp;
float val = x[i];
int64_t tmp_dst = i;
const int64_t i3 = tmp_dst / (ne2*ne1*ne0);
tmp_dst -= i3 * (ne2*ne1*ne0);
const int64_t i2 = tmp_dst / (ne1*ne0);
tmp_dst -= i2 * (ne1*ne0);
const int64_t i1 = tmp_dst / ne0;
tmp_dst -= i1 * ne0;
const int64_t i0 = tmp_dst;
float val = *(const float *) (x + i0*nb00 + i1*nb01 + i2*nb02 + i3*nb03);
if (src1_idx >= 0 && i10 < ne10 && i11 < ne11 && i12 < ne12 && i13 < ne13) {
val += y[((i13*ne12 + i12) * ne11 + i11) * ne10 + i10];
val += *(const float *) (y + i10*nb10 + i11*nb11 + i12*nb12 + i13*nb13);
}
dst[i] = val;
}
@@ -422,15 +434,24 @@ static void gated_op_fused_geglu_quick(const T * x, const T * g, T * dst, const
}
namespace ggml_sycl_detail {
static void acc_f32_sycl(const float *x, const float *y, float *dst,
const int64_t n_elements, const int64_t ne10, const int64_t ne11,
const int64_t ne12, const int64_t ne13, const int64_t s1, const int64_t s2, const int64_t s3,
static void acc_f32_sycl(const char *x, const char *y, float *dst,
const int64_t n_elements,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
const int64_t s1, const int64_t s2, const int64_t s3,
const int64_t offset, queue_ptr stream) {
const int num_blocks = (n_elements + SYCL_ACC_BLOCK_SIZE - 1) / SYCL_ACC_BLOCK_SIZE;
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)),
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
acc_f32(x, y, dst, n_elements, ne10, ne11, ne12, ne13, s1, s2, s3, offset);
acc_f32(x, y, dst, n_elements,
ne0, ne1, ne2, ne3,
nb00, nb01, nb02, nb03,
ne10, ne11, ne12, ne13,
nb10, nb11, nb12, nb13,
s1, s2, s3, offset);
});
}
@@ -843,8 +864,8 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
const float * src0_d = (const float *) src0->data;
const float * src1_d = (const float *) src1->data;
const char * src0_d = (const char *) src0->data;
const char * src1_d = (const char *) src1->data;
float * dst_d = (float *) dst->data;
dpct::queue_ptr stream = ctx.stream();
@@ -853,17 +874,20 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
GGML_ASSERT(src1->type == GGML_TYPE_F32);
GGML_ASSERT( dst->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous(src1));
GGML_ASSERT(dst->nb[0] == ggml_element_size(dst));
GGML_ASSERT(ggml_is_contiguously_allocated(dst));
GGML_ASSERT(ggml_are_same_shape(src0, dst));
const int64_t s1 = dst->op_params[0] / sizeof(float);
const int64_t s2 = dst->op_params[1] / sizeof(float);
const int64_t s3 = dst->op_params[2] / sizeof(float);
const int64_t offset = dst->op_params[3] / sizeof(float);
const int64_t s1 = (int64_t) ((const int32_t *) dst->op_params)[0] / (int64_t) sizeof(float);
const int64_t s2 = (int64_t) ((const int32_t *) dst->op_params)[1] / (int64_t) sizeof(float);
const int64_t s3 = (int64_t) ((const int32_t *) dst->op_params)[2] / (int64_t) sizeof(float);
const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (int64_t) sizeof(float);
ggml_sycl_detail::acc_f32_sycl(src0_d, src1_d, dst_d, ggml_nelements(dst),
dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
s1, s2, s3, offset, stream);
}
+355 -135
View File
@@ -41,7 +41,7 @@
#if SYCL_EXT_ONEAPI_VIRTUAL_MEM
# include <sycl/ext/oneapi/virtual_mem/physical_mem.hpp>
# include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
# define GGML_SYCL_USE_VMM
# define GGML_SYCL_SUPPORT_VMM
#endif
#include <sycl/half_type.hpp>
@@ -74,15 +74,16 @@
#include "ggml-sycl/solve_tri.hpp"
#include "ggml-sycl/gated_delta_net.hpp"
#include "ggml-sycl/pool.hpp"
#include "ggml-sycl/cross_entropy_loss.hpp"
#define MEM_SIZE_2M 0x00200000
#define MEM_SIZE_1G 0x40000000
static bool g_sycl_loaded = false;
int g_ggml_sycl_debug = 0;
int g_ggml_sycl_disable_optimize = 0;
int g_ggml_sycl_disable_graph = 0;
int g_ggml_sycl_disable_dnn = 0;
int g_ggml_sycl_enable_optimize = 1;
int g_ggml_sycl_enable_graph = 0;
int g_ggml_sycl_enable_dnn = 1;
int g_ggml_sycl_enable_vmm = 1;
int g_ggml_sycl_prioritize_dmmv = 0;
int g_ggml_sycl_use_async_mem_op = 0;
@@ -117,7 +118,7 @@ static ggml_sycl_device_info ggml_sycl_init() {
SYCL_CHECK(CHECK_TRY_ERROR(dpct::get_device_info(
prop, device)));
#if !defined(GGML_SYCL_USE_VMM)
#if !defined(GGML_SYCL_SUPPORT_VMM)
info.devices[i].vmm = 0;
#else
info.devices[i].vmm = device.has(sycl::aspect::ext_oneapi_virtual_mem);
@@ -265,14 +266,24 @@ void ggml_backend_sycl_print_sycl_devices() {
print_device_opt_feature(device_count);
}
static const char* dev2dev_int2str(int dev2dev) {
if (dev2dev == DEV2DEV_MEMCPY_SYCL) {
return "SYCL API";
} else if (dev2dev == DEV2DEV_MEMCPY_L0) {
return "Level Zero API";
} else {
return "Unknown";
}
}
static void ggml_check_sycl() try {
static bool initialized = false;
if (!initialized) {
g_ggml_sycl_debug = ggml_sycl_get_env("GGML_SYCL_DEBUG", 0);
g_ggml_sycl_disable_optimize = ggml_sycl_get_env("GGML_SYCL_DISABLE_OPT", 0);
g_ggml_sycl_disable_graph = ggml_sycl_get_env("GGML_SYCL_DISABLE_GRAPH", 1);
g_ggml_sycl_disable_dnn = ggml_sycl_get_env("GGML_SYCL_DISABLE_DNN", 0);
g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1);
g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0);
g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1);
g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1);
g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
@@ -292,66 +303,56 @@ static void ggml_check_sycl() try {
GGML_SYCL_DEBUG("[SYCL] call ggml_check_sycl\n");
GGML_LOG_INFO("Build with Macros:\n");
#if defined(GGML_SYCL_FORCE_MMQ)
GGML_LOG_INFO(" GGML_SYCL_FORCE_MMQ: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_FORCE_MMQ: no\n");
#endif
#if defined(GGML_SYCL_F16)
GGML_LOG_INFO(" GGML_SYCL_F16: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_F16: no\n");
#endif
#if defined(GGML_SYCL_GRAPH)
GGML_LOG_INFO(" GGML_SYCL_GRAPH: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_GRAPH: no\n");
#endif
#if defined(GGML_SYCL_DNNL)
GGML_LOG_INFO(" GGML_SYCL_DNNL: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_DNNL: no\n");
#endif
#if defined(GGML_SYCL_F16)
GGML_LOG_INFO(" GGML_SYCL_F16: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_F16: no\n");
#endif
#if defined(GGML_SYCL_FORCE_MMQ)
GGML_LOG_INFO(" GGML_SYCL_FORCE_MMQ: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_FORCE_MMQ: no\n");
#endif
#if defined(GGML_SYCL_GRAPH)
GGML_LOG_INFO(" GGML_SYCL_GRAPH: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_GRAPH: no\n");
#endif
#if defined(GGML_SYCL_SUPPORT_LEVEL_ZERO_API)
GGML_LOG_INFO(" GGML_SYCL_SUPPORT_LEVEL_ZERO_API: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_SUPPORT_LEVEL_ZERO_API: no\n");
#endif
#if defined(GGML_SYCL_USE_VMM)
GGML_LOG_INFO(" GGML_SYCL_USE_VMM: yes\n");
#if defined(GGML_SYCL_SUPPORT_VMM)
GGML_LOG_INFO(" GGML_SYCL_SUPPORT_VMM: yes\n");
#else
GGML_LOG_INFO(" GGML_SYCL_USE_VMM: no\n");
GGML_LOG_INFO(" GGML_SYCL_SUPPORT_VMM: no\n");
#endif
GGML_LOG_INFO("Running with Environment Variables:\n");
GGML_LOG_INFO(" GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug);
GGML_LOG_INFO(" GGML_SYCL_DISABLE_OPT: %d\n", g_ggml_sycl_disable_optimize);
#ifdef GGML_SYCL_GRAPH
GGML_LOG_INFO(" GGML_SYCL_DISABLE_GRAPH: %d\n", g_ggml_sycl_disable_graph);
#else
GGML_LOG_INFO(" GGML_SYCL_DISABLE_GRAPH: graph disabled by compile flag\n");
#endif
#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
GGML_LOG_INFO(" GGML_SYCL_USE_LEVEL_ZERO_API: %d\n", g_ggml_sycl_use_level_zero_api);
GGML_LOG_INFO(" GGML_SYCL_DEV2DEV_MEMCPY: %d\n", g_ggml_sycl_dev2dev_memcpy);
GGML_LOG_INFO(" GGML_SYCL_DEV2DEV_MEMCPY: %d (%s)\n", g_ggml_sycl_dev2dev_memcpy, dev2dev_int2str(g_ggml_sycl_dev2dev_memcpy));
#else
GGML_LOG_INFO(" GGML_SYCL_USE_LEVEL_ZERO_API: Disable Level Zero API usage by compile flag\n");
GGML_LOG_INFO(" GGML_SYCL_DEV2DEV_MEMCPY: %d, enable to SYCL API since missing GGML_SYCL_SUPPORT_LEVEL_ZERO_API\n",
g_ggml_sycl_dev2dev_memcpy);
GGML_LOG_INFO(" GGML_SYCL_DEV2DEV_MEMCPY: %d (%s), enable to SYCL API since missing GGML_SYCL_SUPPORT_LEVEL_ZERO_API\n",
g_ggml_sycl_dev2dev_memcpy, dev2dev_int2str(g_ggml_sycl_dev2dev_memcpy));
#endif
#if GGML_SYCL_DNNL
GGML_LOG_INFO(" GGML_SYCL_DISABLE_DNN: %d\n", g_ggml_sycl_disable_dnn);
#if defined(GGML_SYCL_DNNL)
GGML_LOG_INFO(" GGML_SYCL_ENABLE_DNN: %d\n", g_ggml_sycl_enable_dnn);
#else
GGML_LOG_INFO(" GGML_SYCL_DISABLE_DNN: DNN disabled by compile flag\n");
GGML_LOG_INFO(" GGML_SYCL_ENABLE_DNN: DNN disabled by compile flag\n");
#endif
#if defined(GGML_SYCL_USE_VMM)
GGML_LOG_INFO(" GGML_SYCL_ENABLE_VMM: %d\n", g_ggml_sycl_enable_vmm);
#else
GGML_LOG_INFO(" GGML_SYCL_ENABLE_VMM: virtual memory extension is not available\n");
#endif
GGML_LOG_INFO(" GGML_SYCL_PRIORITIZE_DMMV: %d\n", g_ggml_sycl_prioritize_dmmv);
g_ggml_sycl_use_async_mem_op_requested = ggml_sycl_get_env("GGML_SYCL_USE_ASYNC_MEM_OP", 1);
GGML_LOG_INFO(" GGML_SYCL_USE_ASYNC_MEM_OP: %d\n", g_ggml_sycl_use_async_mem_op_requested);
#ifdef SYCL_FLASH_ATTN
GGML_LOG_INFO(" GGML_SYCL_ENABLE_FLASH_ATTN: %d\n", g_ggml_sycl_enable_flash_attention);
@@ -360,6 +361,31 @@ static void ggml_check_sycl() try {
g_ggml_sycl_enable_flash_attention);
#endif
#ifdef GGML_SYCL_GRAPH
GGML_LOG_INFO(" GGML_SYCL_ENABLE_GRAPH: %d\n", g_ggml_sycl_enable_graph);
#else
GGML_LOG_INFO(" GGML_SYCL_ENABLE_GRAPH: graph disabled by compile flag\n");
#endif
GGML_LOG_INFO(" GGML_SYCL_ENABLE_OPT: %d\n", g_ggml_sycl_enable_optimize);
#if defined(GGML_SYCL_SUPPORT_VMM)
GGML_LOG_INFO(" GGML_SYCL_ENABLE_VMM: %d\n", g_ggml_sycl_enable_vmm);
#else
GGML_LOG_INFO(" GGML_SYCL_ENABLE_VMM: virtual memory extension is not available\n");
#endif
GGML_LOG_INFO(" GGML_SYCL_PRIORITIZE_DMMV: %d\n", g_ggml_sycl_prioritize_dmmv);
g_ggml_sycl_use_async_mem_op_requested = ggml_sycl_get_env("GGML_SYCL_USE_ASYNC_MEM_OP", 1);
GGML_LOG_INFO(" GGML_SYCL_USE_ASYNC_MEM_OP: %d\n", g_ggml_sycl_use_async_mem_op_requested);
#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
GGML_LOG_INFO(" GGML_SYCL_USE_LEVEL_ZERO_API: %d\n", g_ggml_sycl_use_level_zero_api);
#else
GGML_LOG_INFO(" GGML_SYCL_USE_LEVEL_ZERO_API: Disable Level Zero API usage by compile flag\n");
#endif
GGML_LOG_INFO(" GGML_SYCL_USM_SYSTEM: %d\n", g_ggml_sycl_usm_system);
/* NOT REMOVE, keep it for next optimize for XMX.
@@ -373,7 +399,7 @@ static void ggml_check_sycl() try {
// staging path while preserving queue ordering semantics. Graph support still depends on the extension being
// available, but it no longer needs to control the non-graph fast path.
#if defined(GGML_SYCL_GRAPH) && SYCL_EXT_ONEAPI_ASYNC_MEMORY_ALLOC
g_ggml_sycl_use_async_mem_op = g_ggml_sycl_use_async_mem_op_requested || !g_ggml_sycl_disable_graph;
g_ggml_sycl_use_async_mem_op = g_ggml_sycl_use_async_mem_op_requested || g_ggml_sycl_enable_graph;
if (g_ggml_sycl_use_async_mem_op) {
for (unsigned int i = 0; i < dpct::dev_mgr::instance().device_count(); ++i) {
if (!dpct::dev_mgr::instance().get_device(i).has(sycl::aspect::ext_oneapi_async_memory_alloc)) {
@@ -516,12 +542,14 @@ ggml_backend_sycl_buffer_init_tensor(ggml_backend_buffer_t buffer,
return GGML_STATUS_SUCCESS;
}
if (!g_ggml_sycl_disable_optimize) {
if (g_ggml_sycl_enable_optimize) {
// set reorder extra buffer based on supported type
switch (tensor->type) {
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:{
ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{};
tensor->extra = extra;
@@ -1562,7 +1590,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool {
};
// pool with virtual memory management
#if defined(GGML_SYCL_USE_VMM)
#if defined(GGML_SYCL_SUPPORT_VMM)
struct ggml_sycl_pool_vmm : public ggml_sycl_pool {
static const size_t SYCL_POOL_VMM_MAX_SIZE = 1ull << 35; // 32 GB
@@ -1674,7 +1702,7 @@ struct ggml_sycl_pool_vmm : public ggml_sycl_pool {
GGML_ASSERT(ptr == reinterpret_cast<void *>(pool_addr + pool_used));
}
};
#endif // defined(GGML_SYCL_USE_VMM)
#endif // defined(GGML_SYCL_SUPPORT_VMM)
struct ggml_sycl_pool_host : public ggml_sycl_pool {
queue_ptr qptr;
@@ -1756,11 +1784,11 @@ std::unique_ptr<ggml_sycl_pool> ggml_backend_sycl_context::new_pool_for_host(que
}
std::unique_ptr<ggml_sycl_pool> ggml_backend_sycl_context::new_pool_for_device(queue_ptr qptr, int device) {
#if defined(GGML_SYCL_USE_VMM)
#if defined(GGML_SYCL_SUPPORT_VMM)
if (g_ggml_sycl_enable_vmm && ggml_sycl_info().devices[device].vmm) {
return std::unique_ptr<ggml_sycl_pool>(new ggml_sycl_pool_vmm(qptr, device));
}
#endif // defined(GGML_SYCL_USE_VMM)
#endif // defined(GGML_SYCL_SUPPORT_VMM)
return std::unique_ptr<ggml_sycl_pool>(new ggml_sycl_pool_leg(qptr, device));
}
@@ -2088,11 +2116,148 @@ static int next_power_of_2(int x) {
return n;
}
static void init_argsort_indices_padded(
int * idx,
const int nrows,
const int ncols_pad,
const sycl::nd_item<1> & item_ct1) {
const size_t gid = item_ct1.get_local_range(0) * item_ct1.get_group(0) + item_ct1.get_local_id(0);
const size_t total = (size_t) nrows * (size_t) ncols_pad;
if (gid >= total) {
return;
}
idx[gid] = (int) (gid % (size_t) ncols_pad);
}
template <ggml_sort_order order>
static void argsort_f32_i32_global_pass(const float * x,
int * idx,
const int ncols,
const int nrows,
const int ncols_pad,
const int j,
const int k,
const sycl::nd_item<1> & item_ct1) {
const size_t gid = item_ct1.get_local_range(0) * item_ct1.get_group(0) + item_ct1.get_local_id(0);
const size_t total = (size_t) nrows * (size_t) ncols_pad;
if (gid >= total) {
return;
}
const int row = (int) (gid / (size_t) ncols_pad);
const int col = (int) (gid % (size_t) ncols_pad);
const int ixj = col ^ j;
if (ixj <= col || ixj >= ncols_pad) {
return;
}
const size_t base = (size_t) row * (size_t) ncols_pad;
const size_t pos_a = base + (size_t) col;
const size_t pos_b = base + (size_t) ixj;
const int a = idx[pos_a];
const int b = idx[pos_b];
bool do_swap = false;
if ((col & k) == 0) {
if (a >= ncols ||
(b < ncols &&
(order == GGML_SORT_ORDER_ASC ?
x[(size_t) row * (size_t) ncols + (size_t) a] > x[(size_t) row * (size_t) ncols + (size_t) b] :
x[(size_t) row * (size_t) ncols + (size_t) a] < x[(size_t) row * (size_t) ncols + (size_t) b]))) {
do_swap = true;
}
} else {
if (b >= ncols ||
(a < ncols &&
(order == GGML_SORT_ORDER_ASC ?
x[(size_t) row * (size_t) ncols + (size_t) a] < x[(size_t) row * (size_t) ncols + (size_t) b] :
x[(size_t) row * (size_t) ncols + (size_t) a] > x[(size_t) row * (size_t) ncols + (size_t) b]))) {
do_swap = true;
}
}
if (do_swap) {
idx[pos_a] = b;
idx[pos_b] = a;
}
}
static void copy_argsort_indices_unpadded(const int * idx_padded,
int * dst,
const int nrows,
const int ncols,
const int ncols_pad,
const sycl::nd_item<1> & item_ct1) {
const size_t gid = item_ct1.get_local_range(0) * item_ct1.get_group(0) + item_ct1.get_local_id(0);
const size_t total = (size_t) nrows * (size_t) ncols;
if (gid >= total) {
return;
}
const int row = (int) (gid / (size_t) ncols);
const int col = (int) (gid % (size_t) ncols);
dst[(size_t) row * (size_t) ncols + (size_t) col] = idx_padded[(size_t) row * (size_t) ncols_pad + (size_t) col];
}
static void argsort_f32_i32_sycl(const float *x, int *dst, const int ncols,
const int nrows, ggml_sort_order order,
queue_ptr stream, int device) {
queue_ptr stream, int device, ggml_sycl_pool & pool) {
// bitonic sort requires ncols to be power of 2
const int ncols_pad = next_power_of_2(ncols);
const size_t shared_mem = (size_t) ncols_pad * sizeof(int);
const size_t smpbo = ggml_sycl_info().devices[device].smpbo;
if (shared_mem > smpbo) {
ggml_sycl_pool_alloc<int> idx_padded_alloc(pool, (size_t) nrows * (size_t) ncols_pad);
int * idx_padded = idx_padded_alloc.get();
constexpr size_t block_size = 256;
const size_t total_padded = (size_t) nrows * (size_t) ncols_pad;
const size_t nblocks_padded = (total_padded + block_size - 1) / block_size;
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(nblocks_padded * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item_ct1) { init_argsort_indices_padded(idx_padded, nrows, ncols_pad, item_ct1); });
for (int k = 2; k <= ncols_pad; k *= 2) {
for (int j = k / 2; j > 0; j /= 2) {
if (order == GGML_SORT_ORDER_ASC) {
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(nblocks_padded * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item_ct1) {
argsort_f32_i32_global_pass<GGML_SORT_ORDER_ASC>(x, idx_padded, ncols, nrows, ncols_pad, j,
k, item_ct1);
});
} else if (order == GGML_SORT_ORDER_DESC) {
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(nblocks_padded * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item_ct1) {
argsort_f32_i32_global_pass<GGML_SORT_ORDER_DESC>(x, idx_padded, ncols, nrows, ncols_pad, j,
k, item_ct1);
});
} else {
GGML_ABORT("invalid sort order");
}
}
}
const size_t total = (size_t) nrows * (size_t) ncols;
const size_t nblocks = (total + block_size - 1) / block_size;
stream->parallel_for(sycl::nd_range<1>(sycl::range<1>(nblocks * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item_ct1) {
copy_argsort_indices_unpadded(idx_padded, dst, nrows, ncols, ncols_pad, item_ct1);
});
return;
}
int nth = 1;
int max_block_size = ggml_sycl_info().max_work_group_sizes[device];
@@ -2105,8 +2270,6 @@ static void argsort_f32_i32_sycl(const float *x, int *dst, const int ncols,
const sycl::range<3> block_dims(1, 1, nth);
const sycl::range<3> block_nums(1, nrows, 1);
const size_t shared_mem = ncols_pad * sizeof(int);
GGML_ASSERT(shared_mem<=ggml_sycl_info().devices[device].smpbo);
if (order == GGML_SORT_ORDER_ASC) {
stream->submit([&](sycl::handler &cgh) {
@@ -2429,7 +2592,7 @@ inline void ggml_sycl_op_mul_mat_sycl(
#if GGML_SYCL_DNNL && defined(GGML_SYCL_HAS_BF16)
// Fast path for bf16 src0
if (src0->type == GGML_TYPE_BF16 && !g_ggml_sycl_disable_dnn && ggml_is_contiguous(src0) &&
if (src0->type == GGML_TYPE_BF16 && g_ggml_sycl_enable_dnn && ggml_is_contiguous(src0) &&
row_diff == src0->ne[1]) {
using bf16_t = sycl::ext::oneapi::bfloat16;
ggml_sycl_pool_alloc<bf16_t> src1_as_bf16(ctx.pool(), src1_ncols*ne10);
@@ -2482,7 +2645,7 @@ inline void ggml_sycl_op_mul_mat_sycl(
: src1_as_f16.get();
#if GGML_SYCL_DNNL
if (!g_ggml_sycl_disable_dnn) {
if (g_ggml_sycl_enable_dnn) {
DnnlGemmWrapper::row_gemm(ctx,row_diff, src1_ncols , ne10, src0_ptr,
DnnlGemmWrapper::to_dt<sycl::half>(), src1_ptr, DnnlGemmWrapper::to_dt<sycl::half>(),
dst_dd_i, DnnlGemmWrapper::to_dt<float>(), stream);
@@ -2532,7 +2695,7 @@ inline void ggml_sycl_op_mul_mat_sycl(
const int64_t gemm_flops = (int64_t)row_diff * src1_ncols * ne10;
const bool use_mkl_direct = gemm_flops < 256 * 256 * 256;
#if GGML_SYCL_DNNL
if (!g_ggml_sycl_disable_dnn && !use_mkl_direct) {
if (g_ggml_sycl_enable_dnn && !use_mkl_direct) {
DnnlGemmWrapper::row_gemm(ctx, row_diff, src1_ncols, ne10, src0_ddf_i,
DnnlGemmWrapper::to_dt<float>(), src1_ddf1_i, DnnlGemmWrapper::to_dt<float>(),
dst_dd_i, DnnlGemmWrapper::to_dt<float>(), stream);
@@ -2625,7 +2788,7 @@ inline void ggml_sycl_op_argsort(ggml_backend_sycl_context & ctx, ggml_tensor *
enum ggml_sort_order order = (enum ggml_sort_order) dst->op_params[0];
argsort_f32_i32_sycl(src0_dd, (int *)dst_dd, ncols, nrows, order,
main_stream, ctx.device);
main_stream, ctx.device, ctx.pool());
}
static void ggml_sycl_op_top_k(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
@@ -3352,7 +3515,7 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
const int64_t r3 = ne13 / ne03;
#if GGML_SYCL_DNNL
if (!g_ggml_sycl_disable_dnn) {
if (g_ggml_sycl_enable_dnn) {
int64_t str_a0 = nb00 / type_size_src0;
int64_t str_a1 = nb01 / type_size_src0;
int64_t str_a2 = nb02 / type_size_src0;
@@ -3527,6 +3690,10 @@ inline bool ggml_sycl_supports_reorder_dmmv(enum ggml_type type) {
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
return true;
default:
return false;
@@ -4092,12 +4259,12 @@ static bool reorder_qw(const ggml_tensor * src0, dpct::queue_ptr stream) {
}
static bool should_reorder_tensor(ggml_backend_sycl_context& ctx, const ggml_tensor * dst) {
return !g_ggml_sycl_disable_optimize && //allow optimize, controlled by $GGML_SYCL_DISABLE_OPT
ctx.opt_feature.reorder && //allow this device due to good perf, skip the devices with bad perf.
dst->op == GGML_OP_MUL_MAT && //limit to some supported cases of Q4_0, to do for more cases.
// ne[1] <= 8 so multi-column decode (spec / MTP verify) also bootstraps the reorder;
// all reorderable types have a _switch_ncols kernel.
dst->src[1]->ne[1] <= 8 && dst->src[1]->ne[2]==1 && dst->src[1]->ne[3]==1;
return g_ggml_sycl_enable_optimize && //allow optimize, controlled by $GGML_SYCL_ENABLE_OPT
ctx.opt_feature.reorder && //allow this device due to good perf, skip the devices with bad perf.
dst->op == GGML_OP_MUL_MAT && //limit to some supported cases of Q4_0, to do for more cases.
// ne[1] <= 8 so multi-column decode (spec / MTP verify) also bootstraps the reorder;
// all reorderable types have a _switch_ncols kernel.
dst->src[1]->ne[1] <= 8 && dst->src[1]->ne[2]==1 && dst->src[1]->ne[3]==1;
}
static void opt_for_reorder(ggml_backend_sycl_context * ctx, const ggml_tensor * src0, const ggml_tensor * /* src1 */,
@@ -4136,7 +4303,7 @@ static void opt_for_reorder(ggml_backend_sycl_context * ctx, const ggml_tensor *
// Lazily reorder supported MoE expert weights once their fused path is used.
static void opt_for_reorder_id(ggml_backend_sycl_context * ctx, const ggml_tensor * src0) {
if (g_ggml_sycl_disable_optimize || !ctx->opt_feature.reorder) {
if (!g_ggml_sycl_enable_optimize || !ctx->opt_feature.reorder) {
return;
}
if (src0->type != GGML_TYPE_Q4_K && src0->type != GGML_TYPE_Q5_K && src0->type != GGML_TYPE_Q6_K) {
@@ -4604,6 +4771,11 @@ static void ggml_sycl_im2col_3d(ggml_backend_sycl_context & ctx, ggml_tensor * d
ggml_sycl_op_im2col_3d(ctx, dst);
}
static void ggml_sycl_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
ggml_sycl_op_col2im_1d(ctx, dst);
}
static void ggml_sycl_conv_3d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
ggml_sycl_op_conv_3d(ctx, dst);
@@ -4912,6 +5084,12 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
case GGML_OP_SOFT_MAX_BACK:
ggml_sycl_op_soft_max_back(ctx, dst);
break;
case GGML_OP_CROSS_ENTROPY_LOSS:
ggml_sycl_cross_entropy_loss(ctx, dst);
break;
case GGML_OP_CROSS_ENTROPY_LOSS_BACK:
ggml_sycl_cross_entropy_loss_back(ctx, dst);
break;
case GGML_OP_ROPE:
ggml_sycl_rope(ctx, dst);
break;
@@ -4924,6 +5102,9 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
case GGML_OP_IM2COL_3D:
ggml_sycl_im2col_3d(ctx, dst);
break;
case GGML_OP_COL2IM_1D:
ggml_sycl_col2im_1d(ctx, dst);
break;
case GGML_OP_POOL_2D:
ggml_sycl_pool2d(ctx, dst);
break;
@@ -5204,7 +5385,10 @@ static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_
auto * sycl_ctx = static_cast<ggml_backend_sycl_context *>(backend->context);
#ifdef GGML_SYCL_GRAPH
bool use_sycl_graph = !g_ggml_sycl_disable_graph && check_graph_compatibility(cgraph);
bool use_sycl_graph = false;
if (g_ggml_sycl_enable_graph) {
use_sycl_graph = check_graph_compatibility(cgraph);
}
if (use_sycl_graph) {
const bool graph_support = dpct::get_device(sycl_ctx->device).has(sycl::aspect::ext_oneapi_limited_graph);
if (!graph_support) {
@@ -5470,7 +5654,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
// TODO: This specific configuration can fail with oneDNN and needs more debugging
if (!ggml_is_permuted(a) && ggml_is_permuted(b) && b->ne[2] > 1 && b->ne[3] > 1 &&
a->ne[0] > 128 && a->ne[2] == 1 && src0_type == GGML_TYPE_F16) {
printf("zjy 2\n");
return false;
}
return true;
@@ -5538,70 +5721,99 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
{
ggml_type src0_type = op->src[0]->type;
ggml_type src1_type = op->src[1]->type;
if (src0_type == src1_type && (ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1])) && src0_type != GGML_TYPE_BF16) {
return true;
if (src0_type == GGML_TYPE_F16) {
if (src1_type == GGML_TYPE_Q2_K ||
src1_type == GGML_TYPE_Q3_K ||
src1_type == GGML_TYPE_Q4_K ||
src1_type == GGML_TYPE_Q5_K ||
src1_type == GGML_TYPE_Q6_K ||
src1_type == GGML_TYPE_IQ2_XXS ||
src1_type == GGML_TYPE_IQ2_XS ||
src1_type == GGML_TYPE_IQ2_S ||
src1_type == GGML_TYPE_IQ3_XXS ||
src1_type == GGML_TYPE_IQ1_S ||
src1_type == GGML_TYPE_IQ1_M ||
src1_type == GGML_TYPE_IQ3_S ||
src1_type == GGML_TYPE_IQ4_XS) {
return false;
}
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_F32) {
return true;
if (src0_type == GGML_TYPE_BF16) {
if (src1_type == GGML_TYPE_Q4_0 || //big error in ut
src1_type == GGML_TYPE_Q4_1 || //big error in ut
src1_type == GGML_TYPE_Q8_0 || //big error in ut
src1_type == GGML_TYPE_Q2_K ||
src1_type == GGML_TYPE_Q3_K ||
src1_type == GGML_TYPE_Q4_K ||
src1_type == GGML_TYPE_Q5_K ||
src1_type == GGML_TYPE_Q6_K ||
src1_type == GGML_TYPE_IQ2_XXS ||
src1_type == GGML_TYPE_IQ2_XS ||
src1_type == GGML_TYPE_IQ2_S ||
src1_type == GGML_TYPE_IQ3_XXS ||
src1_type == GGML_TYPE_IQ1_S ||
src1_type == GGML_TYPE_IQ1_M ||
src1_type == GGML_TYPE_IQ3_S ||
src1_type == GGML_TYPE_IQ4_XS) {
return false;
}
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_F16) {
return true;
if (src0_type == GGML_TYPE_F32) {
if (src1_type == GGML_TYPE_Q2_K ||
src1_type == GGML_TYPE_Q3_K ||
src1_type == GGML_TYPE_Q4_K ||
src1_type == GGML_TYPE_Q5_K ||
src1_type == GGML_TYPE_Q6_K ||
src1_type == GGML_TYPE_IQ2_XXS ||
src1_type == GGML_TYPE_IQ2_XS ||
src1_type == GGML_TYPE_IQ2_S ||
src1_type == GGML_TYPE_IQ3_XXS ||
src1_type == GGML_TYPE_IQ1_S ||
src1_type == GGML_TYPE_IQ1_M ||
src1_type == GGML_TYPE_IQ3_S ||
src1_type == GGML_TYPE_IQ4_XS) {
return false;
}
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q8_0) {
return true;
if (src1_type == GGML_TYPE_F32) {
if (src0_type == GGML_TYPE_Q1_0 ||
src0_type == GGML_TYPE_NVFP4 ||
src0_type == GGML_TYPE_Q2_K ||
src0_type == GGML_TYPE_Q3_K ||
src0_type == GGML_TYPE_Q4_K ||
src0_type == GGML_TYPE_Q5_K ||
src0_type == GGML_TYPE_Q6_K ||
src0_type == GGML_TYPE_IQ2_XXS ||
src0_type == GGML_TYPE_IQ2_XS ||
src0_type == GGML_TYPE_IQ2_S ||
src0_type == GGML_TYPE_IQ3_XXS ||
src0_type == GGML_TYPE_IQ1_S ||
src0_type == GGML_TYPE_IQ1_M ||
src0_type == GGML_TYPE_IQ3_S ||
src0_type == GGML_TYPE_IQ4_NL ||
src0_type == GGML_TYPE_IQ4_XS
) {
return false;
}
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q4_0) {
return true;
if (src0_type == src1_type) {
if (src1_type == GGML_TYPE_IQ2_XXS ||
src1_type == GGML_TYPE_IQ2_XS ||
src1_type == GGML_TYPE_IQ2_S ||
src1_type == GGML_TYPE_IQ3_XXS ||
src1_type == GGML_TYPE_IQ3_S ||
src1_type == GGML_TYPE_IQ1_S ||
src1_type == GGML_TYPE_IQ1_M) {
return false;
}
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q4_1) {
return true;
}
if (src0_type == GGML_TYPE_F16 && src1_type == GGML_TYPE_F16) {
return true;
}
if (src0_type == GGML_TYPE_F16 && src1_type == GGML_TYPE_F32) {
return true;
}
if (src0_type == GGML_TYPE_Q8_0 && src1_type == GGML_TYPE_F32) {
return true;
}
if (src0_type == GGML_TYPE_Q4_0 && src1_type == GGML_TYPE_F32) {
return true;
}
if (src0_type == GGML_TYPE_Q4_1 && src1_type == GGML_TYPE_F32) {
return true;
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q5_0) {
return true;
}
if (src0_type == GGML_TYPE_Q5_0 && src1_type == GGML_TYPE_F32) {
return true;
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q5_1) {
return true;
}
if (src0_type == GGML_TYPE_Q5_1 && src1_type == GGML_TYPE_F32) {
return true;
}
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_IQ4_NL) {
return true;
}
if(src0_type == GGML_TYPE_Q8_0 && src1_type == GGML_TYPE_Q8_0) {
return true;
}
if(src0_type == GGML_TYPE_Q5_0 && src1_type == GGML_TYPE_Q5_0) {
return true;
}
if(src0_type == GGML_TYPE_Q5_1 && src1_type == GGML_TYPE_Q5_1) {
return true;
}
if(src0_type == GGML_TYPE_Q4_0 && src1_type == GGML_TYPE_Q4_0) {
return true;
}
if(src0_type == GGML_TYPE_Q4_1 && src1_type == GGML_TYPE_Q4_1) {
return true;
}
return false;
return true;
}
case GGML_OP_REPEAT_BACK:
{
@@ -5643,7 +5855,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
case GGML_OP_SCALE:
return true;
case GGML_OP_CONT:
return op->src[0]->type != GGML_TYPE_BF16;
return true;
case GGML_OP_TRI:
{
const ggml_tensor * src0 = op->src[0];
@@ -5666,6 +5878,14 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
case GGML_OP_IM2COL_3D:
case GGML_OP_UPSCALE:
return true;
case GGML_OP_COL2IM_1D:
return ggml_is_contiguous(op->src[0]) &&
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16
#ifdef GGML_SYCL_HAS_BF16
|| op->type == GGML_TYPE_BF16
#endif
) &&
op->src[0]->type == op->type;
case GGML_OP_CONV_3D:
return op->type == GGML_TYPE_F32 &&
(op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) &&
@@ -5677,8 +5897,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
case GGML_OP_MEAN:
return ggml_is_contiguous(op->src[0]);
case GGML_OP_ARGSORT:
return op->src[0]->ne[0] * sizeof(int) <=
ggml_sycl_info().devices[device].smpbo;
return true;
case GGML_OP_TOP_K: {
const ggml_tensor * src0 = op->src[0];
const int k = op->ne[0];
@@ -5690,9 +5909,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
}
case GGML_OP_POOL_2D:
case GGML_OP_POOL_1D:
return true;
case GGML_OP_ACC:
return ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]);
return true;
case GGML_OP_PAD:
if (ggml_get_op_params_i32(op, 8) != 0) {
return false;
@@ -5725,6 +5943,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
case GGML_OP_FILL:
case GGML_OP_CUMSUM:
case GGML_OP_DIAG:
case GGML_OP_CROSS_ENTROPY_LOSS:
case GGML_OP_CROSS_ENTROPY_LOSS_BACK:
return true;
case GGML_OP_SOLVE_TRI:
return op->src[0]->ne[0] <= SYCL_SOLVE_TRI_MAX_N && op->src[1]->ne[0] <= SYCL_SOLVE_TRI_MAX_K;
+2 -1
View File
@@ -19,6 +19,7 @@
#define WARP_SIZE GGML_SYCL_WARP_SIZE
#define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses
#define SYCL_COL2IM_1D_BLOCK_SIZE 256
#define SYCL_GELU_BLOCK_SIZE 256
#define SYCL_SILU_BLOCK_SIZE 256
#define SYCL_TANH_BLOCK_SIZE 256
@@ -62,7 +63,7 @@
#endif
#ifndef K_QUANTS_PER_ITERATION
#define K_QUANTS_PER_ITERATION 2
#define K_QUANTS_PER_ITERATION 1
#else
static_assert(K_QUANTS_PER_ITERATION == 1 || K_QUANTS_PER_ITERATION == 2, "K_QUANTS_PER_ITERATION must be 1 or 2");
#endif
+18 -15
View File
@@ -129,7 +129,7 @@ typedef struct VkPhysicalDeviceShaderMixedFloatDotProductFeaturesVALVE {
#endif
#define ROUNDUP_POW2(M, N) (((M) + (N) - 1) & ~((N) - 1))
#define CEIL_DIV(M, N) (((M) + (N)-1) / (N))
#define CEIL_DIV(M, N) (((M) / (N)) + (((M) % (N)) != 0))
static bool is_pow2(uint32_t x) { return x > 1 && (x & (x-1)) == 0; }
#define VK_VENDOR_ID_AMD 0x1002
@@ -17370,21 +17370,24 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
return op->type == GGML_TYPE_F32 && op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_SET_ROWS:
{
switch (op->type) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
case GGML_TYPE_BF16:
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
return true;
default:
return false;
if (op->src[0]->type == GGML_TYPE_F32) {
switch (op->type) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
case GGML_TYPE_BF16:
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
return true;
default:
return false;
}
}
return false;
}
case GGML_OP_CONT:
case GGML_OP_CPY:
+11 -1
View File
@@ -525,7 +525,11 @@ const char * ggml_commit(void) {
#if defined(_MSC_VER) || defined(__MINGW32__)
static int64_t timer_freq, timer_start;
void ggml_time_init(void) {
static BOOL CALLBACK ggml_time_init_once(PINIT_ONCE once, PVOID param, PVOID *ctx) {
UNUSED(once);
UNUSED(param);
UNUSED(ctx);
LARGE_INTEGER t;
QueryPerformanceFrequency(&t);
timer_freq = t.QuadPart;
@@ -535,6 +539,12 @@ void ggml_time_init(void) {
// We subtract the program start time to reduce the likelihood of that happening.
QueryPerformanceCounter(&t);
timer_start = t.QuadPart;
return TRUE;
}
void ggml_time_init(void) {
static INIT_ONCE once = INIT_ONCE_STATIC_INIT;
InitOnceExecuteOnce(&once, ggml_time_init_once, NULL, NULL);
}
int64_t ggml_time_ms(void) {
LARGE_INTEGER t;
+8 -2
View File
@@ -186,6 +186,12 @@ function(hf_download version out_var out_resolved)
set(archive "${UI_BINARY_DIR}/dist.tar.gz")
# Use HF_TOKEN to benefit from higher rate limits
set(auth_headers "")
if(DEFINED ENV{HF_TOKEN} AND NOT "$ENV{HF_TOKEN}" STREQUAL "")
list(APPEND auth_headers "HTTPHEADER" "Authorization: Bearer $ENV{HF_TOKEN}")
endif()
set(candidates "")
if(NOT "${version}" STREQUAL "")
list(APPEND candidates "${version}")
@@ -198,7 +204,7 @@ function(hf_download version out_var out_resolved)
message(STATUS "UI: downloading from ${resolved}: ${base}/dist.tar.gz")
file(DOWNLOAD "${base}/dist.tar.gz?download=true" "${archive}"
STATUS status TIMEOUT 300
STATUS status TIMEOUT 300 ${auth_headers}
)
list(GET status 0 rc)
if(NOT rc EQUAL 0)
@@ -208,7 +214,7 @@ function(hf_download version out_var out_resolved)
endif()
file(DOWNLOAD "${base}/dist.tar.gz.sha256?download=true" "${archive}.sha256"
STATUS status TIMEOUT 30
STATUS status TIMEOUT 30 ${auth_headers}
)
list(GET status 0 rc)
if(NOT rc EQUAL 0)
+18 -28
View File
@@ -63,26 +63,6 @@ static bool can_reuse_kq_mask(
// impl
static ggml_tensor * ggml_mul_mat_aux(
ggml_context * ctx,
ggml_tensor * cur,
ggml_tensor * rot) {
const auto n = rot->ne[0];
ggml_tensor * res;
if (!ggml_is_contiguous(cur)) {
res = ggml_cont_2d (ctx, cur, n, ggml_nelements(cur)/n);
} else {
res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
}
res = ggml_mul_mat (ctx, rot, res);
ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
return res;
}
void llm_graph_input_embd::set_input(const llama_ubatch * ubatch) {
if (ubatch->token) {
const int64_t n_tokens = ubatch->n_tokens;
@@ -881,6 +861,14 @@ void llm_graph_input_dsv4::set_input(const llama_ubatch * ubatch) {
dsv4_set_comp_inputs(inp_hca, plan_hca, "hca", debug > 0, ubatch->n_tokens, n_stream);
dsv4_set_comp_inputs(inp_lid, plan_lid, "lid", debug > 0, ubatch->n_tokens, n_stream);
if (inp_csa.k_rot && inp_csa.k_rot->buffer) {
mctx->get_csa()->set_input_k_rot(inp_csa.k_rot);
}
if (inp_hca.k_rot && inp_hca.k_rot->buffer) {
mctx->get_hca()->set_input_k_rot(inp_hca.k_rot);
}
if (inp_lid.k_rot && inp_lid.k_rot->buffer) {
mctx->get_lid()->set_input_k_rot(inp_lid.k_rot);
}
@@ -2633,12 +2621,12 @@ ggml_tensor * llm_graph_context::build_attn(
GGML_ASSERT(v_mla == nullptr);
if (inp->self_k_rot) {
q_cur = ggml_mul_mat_aux(ctx0, q_cur, inp->self_k_rot);
k_cur = ggml_mul_mat_aux(ctx0, k_cur, inp->self_k_rot);
q_cur = llama_mul_mat_hadamard(ctx0, q_cur, inp->self_k_rot);
k_cur = llama_mul_mat_hadamard(ctx0, k_cur, inp->self_k_rot);
}
if (inp->self_v_rot) {
v_cur = ggml_mul_mat_aux(ctx0, v_cur, inp->self_v_rot);
v_cur = llama_mul_mat_hadamard(ctx0, v_cur, inp->self_v_rot);
}
// these nodes are added to the graph together so that they are not reordered
@@ -2669,7 +2657,7 @@ ggml_tensor * llm_graph_context::build_attn(
cb(cur, "kqv_out", il);
if (inp->self_v_rot) {
cur = ggml_mul_mat_aux(ctx0, cur, inp->self_v_rot);
cur = llama_mul_mat_hadamard(ctx0, cur, inp->self_v_rot);
}
if (wo) {
@@ -2874,14 +2862,14 @@ ggml_tensor * llm_graph_context::build_attn(
auto * v_rot = is_swa ? inp->self_v_rot_swa : inp->self_v_rot;
if (k_rot) {
q_cur = ggml_mul_mat_aux(ctx0, q_cur, k_rot);
q_cur = llama_mul_mat_hadamard(ctx0, q_cur, k_rot);
if (k_cur) {
k_cur = ggml_mul_mat_aux(ctx0, k_cur, k_rot);
k_cur = llama_mul_mat_hadamard(ctx0, k_cur, k_rot);
}
}
if (v_rot) {
if (v_cur) {
v_cur = ggml_mul_mat_aux(ctx0, v_cur, v_rot);
v_cur = llama_mul_mat_hadamard(ctx0, v_cur, v_rot);
}
}
@@ -2924,7 +2912,7 @@ ggml_tensor * llm_graph_context::build_attn(
cb(cur, "kqv_out", il);
if (v_rot) {
cur = ggml_mul_mat_aux(ctx0, cur, v_rot);
cur = llama_mul_mat_hadamard(ctx0, cur, v_rot);
}
if (wo) {
@@ -3084,6 +3072,8 @@ llm_graph_input_dsv4 * llm_graph_context::build_inp_dsv4() const {
dsv4_build_comp_inputs(ctx0, inp->inp_csa, mctx_cur->get_csa_plan(ubatch), "csa", n_stream);
dsv4_build_comp_inputs(ctx0, inp->inp_hca, mctx_cur->get_hca_plan(ubatch), "hca", n_stream);
dsv4_build_comp_inputs(ctx0, inp->inp_lid, mctx_cur->get_lid_plan(ubatch), "lid", n_stream);
inp->inp_csa.k_rot = mctx_cur->get_csa()->build_input_k_rot(ctx0);
inp->inp_hca.k_rot = mctx_cur->get_hca()->build_input_k_rot(ctx0);
inp->inp_lid.k_rot = mctx_cur->get_lid()->build_input_k_rot(ctx0);
return (llm_graph_input_dsv4 *) res->add_input(std::move(inp));
+20
View File
@@ -54,6 +54,26 @@ static inline dst_t llama_cast(src_t v) {
}
}
static inline ggml_tensor * llama_mul_mat_hadamard(
ggml_context * ctx,
ggml_tensor * cur,
ggml_tensor * rot) {
const auto n = rot->ne[0];
ggml_tensor * res;
if (!ggml_is_contiguous(cur)) {
res = ggml_cont_2d(ctx, cur, n, ggml_nelements(cur)/n);
} else {
res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
}
res = ggml_mul_mat(ctx, rot, res);
ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
return res;
}
struct time_meas {
time_meas(int64_t & t_acc, bool disable = false);
~time_meas();
+2 -18
View File
@@ -57,22 +57,6 @@ static void ggml_gen_hadamard(ggml_tensor * tensor) {
}
}
static ggml_tensor * ggml_mul_mat_aux(
ggml_context * ctx,
ggml_tensor * cur,
ggml_tensor * rot) {
const auto n = rot->ne[0];
ggml_tensor * res;
res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
res = ggml_mul_mat (ctx, rot, res);
ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
return res;
}
//
// llama_kv_cache
//
@@ -1875,14 +1859,14 @@ ggml_tensor * llama_kv_cache::build_rope_shift(
tmp = ggml_cast(ctx, cur, GGML_TYPE_F32);
// rotate back
tmp = ggml_mul_mat_aux(ctx, tmp, rot);
tmp = llama_mul_mat_hadamard(ctx, tmp, rot);
tmp = ggml_rope_ext(ctx, tmp,
shift, factors, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
yarn_ext_factor, yarn_attn_factor, yarn_beta_fast, yarn_beta_slow);
// rotate fwd
tmp = ggml_mul_mat_aux(ctx, tmp, rot);
tmp = llama_mul_mat_hadamard(ctx, tmp, rot);
tmp = ggml_cpy(ctx, tmp, cur);
} else {
+2
View File
@@ -953,6 +953,8 @@ static buft_list_t make_gpu_buft_list(ggml_backend_dev_t dev, llama_split_mode s
if (buft != nullptr) {
buft_list.emplace_back(dev, buft);
}
} else {
throw std::runtime_error(format("device %s does not support split buffers", ggml_backend_dev_name(dev)));
}
}
+37 -10
View File
@@ -557,7 +557,7 @@ ggml_tensor * llama_model_deepseek4::graph::build_lid_top_k(
cb(indexer_q_pe, "lid_q_pe", il);
indexer_q = ggml_concat(ctx0, indexer_q_nope, indexer_q_pe, 0);
indexer_q = ggml_mul_mat(ctx0, inp_lid.k_rot, indexer_q);
indexer_q = llama_mul_mat_hadamard(ctx0, indexer_q, inp_lid.k_rot);
cb(indexer_q, "lid_q_rot", il);
ggml_tensor * indexer_weights = build_lora_mm(layer.indexer_proj, cur);
@@ -652,10 +652,15 @@ ggml_tensor * llama_model_deepseek4::graph::build_csa_lid_attention(
int il) const {
const auto & inp_csa = inp_dsv4->get_csa();
GGML_ASSERT(inp_csa.kq_mask);
GGML_ASSERT(inp_attn->self_k_rot == nullptr);
ggml_tensor * top_k = build_lid_top_k(model, inp_dsv4, qr, cur, inp_pos, il);
ggml_tensor * k_rot = inp_attn->self_k_rot;
if (k_rot) {
q = llama_mul_mat_hadamard(ctx0, q, k_rot);
kv = llama_mul_mat_hadamard(ctx0, kv, k_rot);
}
ggml_build_forward_expand(gf, q);
ggml_build_forward_expand(gf, kv);
@@ -696,6 +701,9 @@ ggml_tensor * llama_model_deepseek4::graph::build_csa_lid_attention(
ggml_tensor * kq_b = dsv4_build_kq_zero_bias(ctx0, cparams, kq_mask, q->ne[1]);
ggml_tensor * out = build_attn_mha(q, k_all, k_all, kq_b, kq_mask, sinks, nullptr, kq_scale, il);
if (k_rot) {
out = llama_mul_mat_hadamard(ctx0, out, k_rot);
}
cb(out, "attn_csa_lid", il);
return out;
@@ -711,7 +719,12 @@ ggml_tensor * llama_model_deepseek4::graph::build_hca_attention(
int il) const {
const auto & inp_hca = inp_dsv4->get_hca();
GGML_ASSERT(inp_hca.kq_mask);
GGML_ASSERT(inp_attn->self_k_rot == nullptr);
ggml_tensor * k_rot = inp_attn->self_k_rot;
if (k_rot) {
q = llama_mul_mat_hadamard(ctx0, q, k_rot);
kv = llama_mul_mat_hadamard(ctx0, kv, k_rot);
}
ggml_build_forward_expand(gf, q);
ggml_build_forward_expand(gf, kv);
@@ -753,6 +766,9 @@ ggml_tensor * llama_model_deepseek4::graph::build_hca_attention(
ggml_tensor * kq_b = dsv4_build_kq_zero_bias(ctx0, cparams, kq_mask, q->ne[1]);
ggml_tensor * out = build_attn_mha(q, k_all, k_all, kq_b, kq_mask, sinks, nullptr, kq_scale, il);
if (k_rot) {
out = llama_mul_mat_hadamard(ctx0, out, k_rot);
}
cb(out, "attn_hca", il);
return out;
@@ -770,8 +786,8 @@ ggml_tensor * llama_model_deepseek4::graph::build_raw_attention(
ggml_tensor * k_rot = inp_attn->self_k_rot;
if (k_rot) {
q = ggml_mul_mat(ctx0, k_rot, q);
kv = ggml_mul_mat(ctx0, k_rot, kv);
q = llama_mul_mat_hadamard(ctx0, q, k_rot);
kv = llama_mul_mat_hadamard(ctx0, kv, k_rot);
}
ggml_build_forward_expand(gf, q);
@@ -788,6 +804,9 @@ ggml_tensor * llama_model_deepseek4::graph::build_raw_attention(
ggml_tensor * kq_b = dsv4_build_kq_zero_bias(ctx0, cparams, kq_mask, q->ne[1]);
ggml_tensor * out = build_attn_mha(q, k, k, kq_b, kq_mask, sinks, nullptr, kq_scale, il);
if (k_rot) {
out = llama_mul_mat_hadamard(ctx0, out, k_rot);
}
cb(out, "attn_raw", il);
return out;
@@ -917,6 +936,11 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
"csa_state_compress",
il);
if (inp_dsv4->get_csa().k_rot) {
kv_comp_csa_state = llama_mul_mat_hadamard(ctx0, kv_comp_csa_state, inp_dsv4->get_csa().k_rot);
cb(kv_comp_csa_state, "csa_state_compress_rot", il);
}
ggml_build_forward_expand(gf, inp_dsv4->mctx->get_csa()->cpy_k(ctx0,
kv_comp_csa_state, inp_dsv4->get_csa().state_write_idxs, il));
@@ -965,7 +989,7 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
il);
if (inp_dsv4->get_lid().k_rot) {
kv_comp_lid_state = ggml_mul_mat(ctx0, inp_dsv4->get_lid().k_rot, kv_comp_lid_state);
kv_comp_lid_state = llama_mul_mat_hadamard(ctx0, kv_comp_lid_state, inp_dsv4->get_lid().k_rot);
cb(kv_comp_lid_state, "lid_state_compress_rot", il);
}
@@ -1007,6 +1031,11 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
"hca_state_compress",
il);
if (inp_dsv4->get_hca().k_rot) {
kv_comp_hca = llama_mul_mat_hadamard(ctx0, kv_comp_hca, inp_dsv4->get_hca().k_rot);
cb(kv_comp_hca, "hca_state_compress_rot", il);
}
ggml_build_forward_expand(gf, inp_dsv4->mctx->get_hca()->cpy_k(ctx0,
kv_comp_hca, inp_dsv4->get_hca().state_write_idxs, il));
hca_state_dep = kv_comp_hca;
@@ -1035,13 +1064,11 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
if (ratio == DSV4_CSA_RATIO &&
inp_dsv4->get_csa().kq_mask &&
inp_dsv4->get_lid().kq_mask &&
inp_dsv4->get_lid().k_rot &&
inp_attn->self_k_rot == nullptr) {
inp_dsv4->get_lid().k_rot) {
out = build_csa_lid_attention(model, inp_dsv4, inp_attn, q, kv, qr, cur, inp_pos, layer.attn_sinks,
1.0f/sqrtf(float(n_embd_head)), il);
} else if (ratio == DSV4_HCA_RATIO &&
inp_dsv4->get_hca().kq_mask &&
inp_attn->self_k_rot == nullptr) {
inp_dsv4->get_hca().kq_mask) {
out = build_hca_attention(inp_dsv4, inp_attn, q, kv, layer.attn_sinks,
1.0f/sqrtf(float(n_embd_head)), il);
} else {
+10 -4
View File
@@ -523,6 +523,7 @@ void server_models::load_models() {
// collect all threads to join in one pass while the lock is held:
// - monitoring threads from just-unloaded models (to_unload)
// - threads of finished downloads (DOWNLOADED), they acquire the mutex on exit
// - threads of already-UNLOADED models that are being removed from source
std::vector<std::thread> threads_to_join;
for (const auto & name : to_unload) {
@@ -535,6 +536,13 @@ void server_models::load_models() {
if (inst.meta.status == SERVER_MODEL_STATUS_DOWNLOADING) {
continue; // downloading models are not from config sources, leave them alone
}
if (inst.meta.status == SERVER_MODEL_STATUS_DOWNLOADED) {
// joining this thread under the lock deadlocks: it locks the mutex on its way out
if (inst.th.joinable()) {
threads_to_join.push_back(std::move(inst.th));
}
continue;
}
if (final_presets.find(name) == final_presets.end() && !inst.meta.is_running() && inst.th.joinable()) {
threads_to_join.push_back(std::move(inst.th));
}
@@ -550,10 +558,8 @@ void server_models::load_models() {
if (it->second.meta.status == SERVER_MODEL_STATUS_DOWNLOADING) {
++it; // download thread is still busy, skip
} else if (it->second.meta.status == SERVER_MODEL_STATUS_DOWNLOADED) {
// download finished, safe to erase
if (it->second.th.joinable()) {
it->second.th.join();
}
// download finished, thread is joined above, safe to erase
GGML_ASSERT(!it->second.th.joinable());
it = mapping.erase(it);
} else if (final_presets.find(it->first) == final_presets.end()) {
SRV_INF("(reload) removing model name=%s (no longer in source)\n", it->first.c_str());
+5 -2
View File
@@ -31,6 +31,9 @@ import wget
DEFAULT_HTTP_TIMEOUT = 60
# per-request timeout, a hung server fails the test instead of stalling the CI for hours
DEFAULT_REQUEST_TIMEOUT = 600
class ServerResponse:
headers: dict
@@ -330,7 +333,7 @@ class ServerProcess:
path: str,
data: dict | Any | None = None,
headers: dict | None = None,
timeout: float | None = None,
timeout: float | None = DEFAULT_REQUEST_TIMEOUT,
) -> ServerResponse:
url = f"http://{self.server_host}:{self.server_port}{path}"
parse_body = False
@@ -389,7 +392,7 @@ class ServerProcess:
path: str,
data: dict | None = None,
headers: dict | None = None,
timeout: float | None = None,
timeout: float | None = DEFAULT_REQUEST_TIMEOUT,
) -> dict:
stream = data.get('stream', False)
if stream:
@@ -27,7 +27,10 @@
let { onSearchClick = () => {} }: Props = $props();
const { handleKeydown } = useKeyboardShortcuts({ activateSearchMode: () => onSearchClick() });
const { handleKeydown } = useKeyboardShortcuts({
activateSearchMode: () => onSearchClick(),
toggleSidebar: () => toggleExpandedMode()
});
let isExpandedMode = $state(false);
let hoveredTooltip = $state<string | null>(null);
+1
View File
@@ -9,6 +9,7 @@ export enum KeyboardKey {
ARROW_LEFT = 'ArrowLeft',
ARROW_RIGHT = 'ArrowRight',
TAB = 'Tab',
B_LOWER = 'b',
D_LOWER = 'd',
D_UPPER = 'D',
E_UPPER = 'E',
@@ -9,6 +9,7 @@ interface KeyboardShortcutsCallbacks {
deleteActiveConversation?: () => void;
navigateToPrevConversation?: () => void;
navigateToNextConversation?: () => void;
toggleSidebar?: () => void;
}
export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
@@ -21,6 +22,11 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
callbacks.onSearchActivated?.();
}
if (isCmdOrCtrl && event.key === KeyboardKey.B_LOWER) {
event.preventDefault();
callbacks.toggleSidebar?.();
}
if (
isCmdOrCtrl &&
event.shiftKey &&
+24
View File
@@ -314,6 +314,30 @@ export class MCPService {
)
);
if (method === 'DELETE' && url.includes(CORS_PROXY_ENDPOINT)) {
const response = new Response(null, { status: 200, statusText: 'OK' });
logIfEnabled(
this.createLog(
MCPConnectionPhase.INITIALIZING,
`HTTP 200 ${method} ${url} (fake response)`,
MCPLogLevel.INFO,
{
response: {
url,
status: response.status,
statusText: response.statusText,
durationMs: 0,
isFake: true
}
}
)
);
// fake response, bypass real fetch()
return response;
}
try {
const response = await fetch(input, {
...baseInit,
+26
View File
@@ -154,6 +154,32 @@ describe('MCPService', () => {
});
});
it('DELETE request with CORS proxy should return a fake 200 response', async () => {
const logs: MCPConnectionLog[] = [];
const fetchMock = vi.fn();
vi.stubGlobal('fetch', fetchMock);
const config: MCPServerConfig = {
url: 'https://example.com/mcp',
transport: MCPTransportType.STREAMABLE_HTTP,
useProxy: true
};
const controller = createDiagnosticFetch(config, (log) => logs.push(log), {}, true);
const response = await controller.fetch(
'http://localhost:8080/cors-proxy?url=https%3A%2F%2Fexample.com%2Fmcp',
{ method: 'DELETE' }
);
expect(fetchMock).not.toHaveBeenCalled();
expect(response.status).toBe(200);
expect(logs.at(-1)?.details).toMatchObject({
response: { status: 200, isFake: true }
});
});
it('partially redacts mcp-session-id in diagnostic request and response logs', async () => {
const logs: MCPConnectionLog[] = [];
const response = new Response('{}', {