alu@king
OS Garuda Linux x86_64
├ Kernel Linux 6.17.9-zen1-1-zen
├ Packages 1557 (pacman)[stable]
└ Shell fish 4.2.1
DE KDE Plasma 6.5.3
├ Window Manager KWin (Wayland)
├ Login Manager sddm-autologin 0.21.0 (Wayland)
├ WM Theme Sweet-Dark
├ Color Themes Dr460nized (Sweet) [Qt], Sweet-Dark [GTK2/3/4]
├ System Icons breeze-dark [Qt], breeze-dark [GTK2/3/4]
├ System Fonts Fira Sans Heavy (10pt, Regular) [Qt], Fira Sans Heavy (10pt, ExtraLight) [GTK2/3/4]
└ Terminal konsole 25.8.3
PC Notebook (3.0)
├ CPU 12th Gen Intel(R) Core(TM) i7-12700H (20) @ 3.50 GHz
├ GPU Intel Arc A730M @ 2.05 GHz [Discrete]
├ GPU Intel Iris Xe Graphics @ 1.40 GHz [Integrated]
├ Vulkan 1.4.318 - Intel open-source Mesa driver [Mesa 25.2.7-arch1.1]
└ Display(s) 1920x1080 @ 1.13x in 16", 144 Hz [Built-in]
^[[A
:: initializing oneAPI environment ...
bash: BASH_VERSION = 5.3.3(1)-release
args: Using "$@" for setvars.sh arguments: --force
:: advisor -- latest
:: ccl -- latest
:: compiler -- latest
:: dal -- latest
:: debugger -- latest
:: dev-utilities -- latest
:: dnnl -- latest
:: dpcpp-ct -- latest
:: dpl -- latest
:: ipp -- latest
:: ippcp -- latest
:: mkl -- latest
:: mpi -- latest
:: pti -- latest
:: tbb -- latest
:: umf -- latest
:: vtune -- latest
:: oneAPI environment initialized ::
╭─alu@king in ~
╰─λ ./XAIGPUARC.sh
✅ ✅ALLE PATCHES ERFOLGREICH ANGEWAND
🔷 HOLE ONE API KOEPF
🔷 SETVARS.SH SETZEN UND🔍
:: initializing oneAPI environment ...
XAIGPUARC.sh: BASH_VERSION = 5.3.3(1)-release
args: Using "$@" for setvars.sh arguments: --force
:: advisor -- latest
:: ccl -- latest
:: compiler -- latest
:: dal -- latest
:: debugger -- latest
:: dev-utilities -- latest
:: dnnl -- latest
:: dpcpp-ct -- latest
:: dpl -- latest
:: ipp -- latest
:: ippcp -- latest
:: mkl -- latest
:: mpi -- latest
:: pti -- latest
:: tbb -- latest
:: umf -- latest
:: vtune -- latest
:: oneAPI environment initialized ::
🔷 ✅ VERBINDUNG ONEAPI GELADEN... (DPCPP_ROOT=/opt/intel/oneapi/compiler/2025.0 und MKL_ROOT=/opt/intel/oneapi/mkl/2025.0)
✅ ✅GEFUNDENE-AKTUELLE XAIGPUARC VERSION-NEUBAU-UNNÖTIG-FORTFAHREN./XAIGPUARC/bin/llama-cli und ./XAIGPUARC/bin/llama-ls-sycl-device
🔷 ✅->ÜBERSPRINGE-BAUVORGANG
🔷 ⚙->UPDATE-JETZT-NEUESTE-LLAMA-VERSION-BITTE-WARTEN
🔷 📦 BAUE-XAIGPUARC-BITTE WARTEN
🔷 🔍->AKTUALISIERE UNTERMODULE
remote: Enumerating objects: 31, done.
remote: Counting objects: 100% (19/19), done.
remote: Compressing objects: 100% (17/17), done.
remote: Total 31 (delta 6), reused 5 (delta 2), pack-reused 12 (from 3)
Entpacke Objekte: 100% (31/31), 108.70 KiB | 1.98 MiB/s, fertig.
Von https://github.com/ggerganov/llama.cpp
3659aa28e..9d0229967 master -> origin/master
- [neues Tag] b7266 -> b7266
- [neues Tag] b7268 -> b7268
- [neues Tag] b7270 -> b7270
Aktualisiere 3659aa28e..9d0229967
Fast-forward
common/arg.cpp | 2 +-
examples/simple-cmake-pkg/README.md | 1 +
ggml/src/ggml-cpu/arch/arm/repack.cpp | 2 --
tools/server/server-context.cpp | 171 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++--------------------------------------------------------------
tools/server/server-models.cpp | 1 +
tools/server/server-queue.cpp | 10 ++++++++
tools/server/server-queue.h | 9 +++++++
tools/server/server-task.cpp | 27 +++++++++++++++++---
tools/server/server-task.h | 40 ++++++++++++++++++++++++++---
9 files changed, 168 insertions(+), 95 deletions(-)
✅ ✅ LLAMA.CPP ANTWORTET..UNTERGRUPPEN WERDEN GELADEN
🔷 🔷 🏗 🩹 Patches für ggml-sycl anwenden (Header & CMake & Kernel-Dispatch-Registrierung)
🔷 🔷->PATCH 1/6: DOCTPHELPER FEHLGESCHLAGEN. ABHÄNGIGKEITSLISTE PRÜFEN
🔷 🔷-> ✅ PATCH 1/6 ERFOLGREICH
🔷 🔷->PATCH 2/6: ggml_flash_attention_sycl
🔷 🔷-> CMakeLists.txt für Kernel als OBJECT-Library erstellt
🔷 🔷->✅🏗PATCH 2/6 ERFOLGREICH ggml_flash_attention_sycl zu Haupt-CMake hinzugefügt
🔷 🔷-> PATCH 3/6: CMakeLists.txt anpassen (Alle Header-Pfade für icpx).
🔷 🔷->✅🏗PATCH 3/6 erfolgreich: Alle Header-Pfade injiziert.
🔷 🔷->🏗PATCH 4/6: ggml_flash_attention_sycl.cpp INJIZIEREN🏗
🔷 🔷->PATCH 4/6 DEKLARATION ERFOLGREICH EINGEFÜGT
🔷 🔷->Versuche, den Dispatch-Case (FA) mittels AWK einzufügen.
🔷 🔷->PATCH 4/6 ERFOLGREICH✅UNTERBAU ERFOLGREICH EINGEFÜHRT✅
🔷 🔷->✅PATCH 4/6 ERFOLGREICH-FLASHATTENTENTION-GELADEN
🔷 🔷->PATCH 5/6: INJIZIEREN OBJEKT🏗VARIABLEN AUS UNTERBLOCK VON SYCL BIBLIOTHEKEN..
🔷 🔷->5a/6: OBJEKT VARIABLEN 🏗 ERFOLGREICH DEFINIERT
🔷 🔷->⚠PATCH 5b/6 IST BEREITS AKTIV INJECTION WIRD ÜBERSPRUNGEN
🔷 🔷->PATCH 6/6: ssm_conv.cpp WARNUNG beheben VORZEICHENVERGLEICH
🔷 🔷->⚠PATCH 6/6ssm_conv.cppZEILE-NICHT-GEFUNDEN-UEBERSPRINGE
🔷 🔍SUCHE NACH VERFÜGBAREN SYCL GERÄTEN AUF IHREM SYSTEM.
Found 2 SYCL devices:
| | | | |Max | |Max |Global | |
| | | | |compute|Max work|sub |mem | |
|ID| Device Type| Name|Version|units |group |group|size | Driver version|
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|
| 0| [level_zero:gpu:0]| Intel Arc A730M Graphics| 12.55| 384| 1024| 32| 12160M| 1.13.36015|
| 1| [level_zero:gpu:1]| Intel Iris Xe Graphics| 12.3| 96| 512| 32| 30705M| 1.13.36015|
SYCL Optimization Feature:
|ID| Device Type|Reorder|
|--|-------------------|-------|
| 0| [level_zero:gpu:0]| Y|
| 1| [level_zero:gpu:1]| Y|
⚠ ⚠KEINE KOMPATIBLEN SYCL GERÄTE GEFUNDEN: ERROR❌AKTUELLE ABHÄNGIGKEITEN PRÜFEN
🔷 🔍SUCHE SYCL FÄHIGES GERÄT 🏗 AUF IHREM SYSTEM
Found 1 SYCL devices:
| | | | |Max | |Max |Global | |
| | | | |compute|Max work|sub |mem | |
|ID| Device Type| Name|Version|units |group |group|size | Driver version|
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|
| 0| [level_zero:gpu:0]| Intel Arc A730M Graphics| 12.55| 384| 1024| 32| 12160M| 1.13.36015|
SYCL Optimization Feature:
|ID| Device Type|Reorder|
|--|-------------------|-------|
| 0| [level_zero:gpu:0]| Y|
🔷 🚀STARTE KI ANTWORT PER F16 INFERENCE AUF IHRER iGPU/dGPU MIT FOLGENDEN PARAMETERNARC (ID: 0->❌ANBINDUNG FEHLGESCHLAGEN) with ngl=0 using ./XAIGPUARC/bin/llama-cli...
build: 7254 (e7c2cf135) with Intel(R) oneAPI DPC/C Compiler 2025.0.4 (2025.0.4.20241205) for x86_64-unknown-linux-gnu
main: llama backend init
main: load the model and apply lora adapter, if any
llama_model_load_from_file_impl: using device SYCL0 (Intel(R) Arc(TM) A730M Graphics) (unknown id) - 11597 MiB free
llama_model_loader: loaded meta data with 48 key-value pairs and 309 tensors from models/Ministral-3-8B-Instruct-2512-UD-Q8_K_XL.gguf (version GGUF V3 (latest))
llama_model_loader: Dumping metadata keys/values. Note: KV overrides do not apply in this output.
llama_model_loader: - kv 0: general.architecture str = mistral3
llama_model_loader: - kv 1: general.type str = model
llama_model_loader: - kv 2: general.name str = Ministral-3-8B-Instruct-2512
llama_model_loader: - kv 3: general.version str = 2512
llama_model_loader: - kv 4: general.finetune str = Instruct
llama_model_loader: - kv 5: general.basename str = Ministral-3-8B-Instruct-2512
llama_model_loader: - kv 6: general.quantized_by str = Unsloth
llama_model_loader: - kv 7: general.size_label str = 8B
llama_model_loader: - kv 8: general.repo_url str = https://huggingface.co/unsloth
llama_model_loader: - kv 9: mistral3.block_count u32 = 34
llama_model_loader: - kv 10: mistral3.context_length u32 = 262144
llama_model_loader: - kv 11: mistral3.embedding_length u32 = 4096
llama_model_loader: - kv 12: mistral3.feed_forward_length u32 = 14336
llama_model_loader: - kv 13: mistral3.attention.head_count u32 = 32
llama_model_loader: - kv 14: mistral3.attention.head_count_kv u32 = 8
llama_model_loader: - kv 15: mistral3.attention.layer_norm_rms_epsilon f32 = 0,000010
llama_model_loader: - kv 16: mistral3.attention.key_length u32 = 128
llama_model_loader: - kv 17: mistral3.attention.value_length u32 = 128
llama_model_loader: - kv 18: mistral3.vocab_size u32 = 131072
llama_model_loader: - kv 19: mistral3.rope.dimension_count u32 = 128
llama_model_loader: - kv 20: mistral3.rope.scaling.type str = yarn
llama_model_loader: - kv 21: mistral3.rope.scaling.factor f32 = 16,000000
llama_model_loader: - kv 22: mistral3.rope.scaling.yarn_beta_fast f32 = 32,000000
llama_model_loader: - kv 23: mistral3.rope.scaling.yarn_beta_slow f32 = 1,000000
llama_model_loader: - kv 24: mistral3.rope.scaling.yarn_log_multiplier f32 = 1,000000
llama_model_loader: - kv 25: mistral3.rope.scaling.original_context_length u32 = 16384
llama_model_loader: - kv 26: mistral3.rope.freq_base f32 = 1000000,000000
llama_model_loader: - kv 27: mistral3.attention.temperature_scale f32 = 0,100000
llama_model_loader: - kv 28: tokenizer.ggml.model str = gpt2
llama_model_loader: - kv 29: tokenizer.ggml.pre str = tekken
llama_model_loader: - kv 30: tokenizer.ggml.tokens arr[str,131072] = ["", "", "", "[INST]", "[...
llama_model_loader: - kv 31: tokenizer.ggml.token_type arr[i32,131072] = [3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, ...
llama_model_loader: - kv 32: tokenizer.ggml.merges arr[str,269443] = ["Ġ Ġ", "Ġ t", "e r", "i n", "Ġ �...
llama_model_loader: - kv 33: tokenizer.ggml.bos_token_id u32 = 1
llama_model_loader: - kv 34: tokenizer.ggml.eos_token_id u32 = 2
llama_model_loader: - kv 35: tokenizer.ggml.unknown_token_id u32 = 0
llama_model_loader: - kv 36: tokenizer.ggml.padding_token_id u32 = 11
llama_model_loader: - kv 37: tokenizer.ggml.add_bos_token bool = true
llama_model_loader: - kv 38: tokenizer.ggml.add_sep_token bool = false
llama_model_loader: - kv 39: tokenizer.ggml.add_eos_token bool = false
llama_model_loader: - kv 40: tokenizer.chat_template str = {#- Unsloth template fixes #}\n{#- Def...
llama_model_loader: - kv 41: tokenizer.ggml.add_space_prefix bool = false
llama_model_loader: - kv 42: general.quantization_version u32 = 2
llama_model_loader: - kv 43: general.file_type u32 = 7
llama_model_loader: - kv 44: quantize.imatrix.file str = Ministral-3-8B-Instruct-2512-GGUF/ima...
llama_model_loader: - kv 45: quantize.imatrix.dataset str = unsloth_calibration_Ministral-3-8B-In...
llama_model_loader: - kv 46: quantize.imatrix.entries_count u32 = 238
llama_model_loader: - kv 47: quantize.imatrix.chunks_count u32 = 82
llama_model_loader: - type f32: 69 tensors
llama_model_loader: - type f16: 61 tensors
llama_model_loader: - type q8_0: 179 tensors
print_info: file format = GGUF V3 (latest)
print_info: file type = Q8_0
print_info: file size = 10,32 GiB (10,45 BPW)
load: special_eos_id is not in special_eog_ids - the tokenizer config may be incorrect
load: printing all EOG tokens:
load: - 2 ('')
load: special tokens cache size = 1000
load: token to piece cache size = 0,8498 MB
print_info: arch = mistral3
print_info: vocab_only = 0
print_info: n_ctx_train = 262144
print_info: n_embd = 4096
print_info: n_embd_inp = 4096
print_info: n_layer = 34
print_info: n_head = 32
print_info: n_head_kv = 8
print_info: n_rot = 128
print_info: n_swa = 0
print_info: is_swa_any = 0
print_info: n_embd_head_k = 128
print_info: n_embd_head_v = 128
print_info: n_gqa = 4
print_info: n_embd_k_gqa = 1024
print_info: n_embd_v_gqa = 1024
print_info: f_norm_eps = 0,0e+00
print_info: f_norm_rms_eps = 1,0e-05
print_info: f_clamp_kqv = 0,0e+00
print_info: f_max_alibi_bias = 0,0e+00
print_info: f_logit_scale = 0,0e+00
print_info: f_attn_scale = 0,0e+00
print_info: n_ff = 14336
print_info: n_expert = 0
print_info: n_expert_used = 0
print_info: n_expert_groups = 0
print_info: n_group_used = 0
print_info: causal attn = 1
print_info: pooling type = 0
print_info: rope type = 0
print_info: rope scaling = yarn
print_info: freq_base_train = 1000000,0
print_info: freq_scale_train = 0,0625
print_info: n_ctx_orig_yarn = 16384
print_info: rope_finetuned = unknown
print_info: model type = 8B
print_info: model params = 8,49 B
print_info: general.name = Ministral-3-8B-Instruct-2512
print_info: vocab type = BPE
print_info: n_vocab = 131072
print_info: n_merges = 269443
print_info: BOS token = 1 ''
print_info: EOS token = 2 ''
print_info: UNK token = 0 ''
print_info: PAD token = 11 ''
print_info: LF token = 1010 'Ċ'
print_info: EOG token = 2 ''
print_info: max token length = 150
load_tensors: loading model tensors, this can take a while... (mmap = true)
load_tensors: offloading 34 repeating layers to GPU
load_tensors: offloading output layer to GPU
load_tensors: offloaded 35/35 layers to GPU
load_tensors: CPU_Mapped model buffer size = 1024,00 MiB
load_tensors: SYCL0 model buffer size = 9547,83 MiB
................................................................................
llama_context: constructing llama_context
llama_context: n_seq_max = 1
llama_context: n_ctx = 2048
llama_context: n_ctx_seq = 2048
llama_context: n_batch = 1024
llama_context: n_ubatch = 512
llama_context: causal_attn = 1
llama_context: flash_attn = auto
llama_context: kv_unified = false
llama_context: freq_base = 1000000,0
llama_context: freq_scale = 0,0625
llama_context: n_ctx_seq (2048) < n_ctx_train (262144) -- the full capacity of the model will not be utilized
Running with Environment Variables:
GGML_SYCL_DEBUG: 0
GGML_SYCL_DISABLE_OPT: 0
GGML_SYCL_DISABLE_GRAPH: 1
GGML_SYCL_DISABLE_DNN: 0
GGML_SYCL_PRIORITIZE_DMMV: 0
Build with Macros:
GGML_SYCL_FORCE_MMQ: no
GGML_SYCL_F16: yes
Found 1 SYCL devices:
| | | | |Max | |Max |Global | |
| | | | |compute|Max work|sub |mem | |
|ID| Device Type| Name|Version|units |group |group|size | Driver version|
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|
| 0| [level_zero:gpu:0]| Intel Arc A730M Graphics| 12.55| 384| 1024| 32| 12160M| 1.13.36015|
SYCL Optimization Feature:
|ID| Device Type|Reorder|
|--|-------------------|-------|
| 0| [level_zero:gpu:0]| Y|
llama_context: SYCL_Host output buffer size = 0,50 MiB
llama_kv_cache: SYCL0 KV buffer size = 272,00 MiB
llama_kv_cache: size = 272,00 MiB ( 2048 cells, 34 layers, 1/1 seqs), K (f16): 136,00 MiB, V (f16): 136,00 MiB
llama_context: layer 0 is assigned to device SYCL0 but the Flash Attention tensor is assigned to device CPU (usually due to missing support)
llama_context: Flash Attention was auto, set to disabled
llama_context: SYCL0 compute buffer size = 264,00 MiB
llama_context: SYCL_Host compute buffer size = 12,02 MiB
llama_context: graph nodes = 1264
llama_context: graph splits = 2
common_init_from_params: added logit bias = -inf
common_init_from_params: setting dry_penalty_last_n to ctx_size = 2048
common_init_from_params: warming up the model with an empty run - please wait ... (--no-warmup to disable)
main: llama threadpool init, n_threads = 4
system_info: n_threads = 4 (n_threads_batch = 4) / 20 | CPU : SSE3 = 1 | SSSE3 = 1 | AVX = 1 | AVX_VNNI = 1 | AVX2 = 1 | F16C = 1 | FMA = 1 | BMI2 = 1 | LLAMAFILE = 1 | OPENMP = 1 | REPACK = 1 |
sampler seed: 1530963926
sampler params:
repeat_last_n = 64, repeat_penalty = 1,000, frequency_penalty = 0,000, presence_penalty = 0,000
dry_multiplier = 0,000, dry_base = 1,750, dry_allowed_length = 2, dry_penalty_last_n = 2048
top_k = 40, top_p = 0,950, min_p = 0,050, xtc_probability = 0,000, xtc_threshold = 0,100, typical_p = 1,000, top_n_sigma = -1,000, temp = 0,800
mirostat = 0, mirostat_lr = 0,100, mirostat_ent = 5,000
sampler chain: logits -> logit-bias -> penalties -> dry -> top-n-sigma -> top-k -> typical -> top-p -> min-p -> xtc -> temp-ext -> dist
generate: n_ctx = 2048, n_batch = 1024, n_predict = 1024, n_keep = 1
IMPORTANT: The current llama-cli will be moved to llama-completion in the near future
New llama-cli will have enhanced features and improved user experience
More info: https://github.com/ggml-org/llama.cpp/discussions/17618
//SYCL Flash Attention Kernel für llama.cpp
// Optimierte Implementation mit Tiling-Strategie und Vektorisierung
#include <sycl/sycl.hpp>
#include <sycl/ext/intel/math.hpp>
#include
#include
#include ggml-sycl.h
#include ggml-impl.h
using namespace sycl;
// Kernel-Parameter (optimiert für Intel Arc GPUs)
constexpr int BLOCK_M = 64; // Queries pro Block
constexpr int BLOCK_N = 128; // Keys pro Block (Tiling-Größe)
constexpr int D_MAX = 128; // Maximale Head-Dimension
constexpr int VEC_SIZE = 8; // Vektorgröße für Intel Arc (8 für optimalen SIMD)
// Vektorisierte Hilfsfunktionen
/**
Vektorisiertes Dot-Produkt zwischen Q[i] und K[j]
scalar_t Datentyp (sycl::half oder float)
Q_row_float Query-Zeile als float[]
Dot-Product als float
/
template
float dot_product_vec(const float Q_row_float, const scalar_t* K_ptr, int d_k) {
if constexpr (std::is_same_v<scalar_t, sycl::half>) {
if (d_k % VEC_SIZE != 0) {
// Fallback für nicht-vektorisierte Dimensionen
float score = 0.0f;
for (int di = 0; di < d_k; ++di) {
score += Q_row_float[di] * static_cast(K_ptr[di]);
}
return score;
}
constexpr int vec_elements = VEC_SIZE;
using vec_half = sycl::vec<sycl::half, vec_elements>;
using vec_float = sycl::vec<float, vec_elements>;
float final_score = 0.0f;
int vec_iters = d_k / vec_elements;
for (int v = 0; v < vec_iters; ++v) {
// Lade K-Vektor (half)
vec_half k_half_vec;
k_half_vec.load(v * vec_elements, K_ptr);
// Konvertiere zu float
vec_float k_float_vec = k_half_vec.template convert();
// Lade Q-Vektor (float)
vec_float q_float_vec;
q_float_vec.load(v * vec_elements, Q_row_float);
// Vektorisiertes Dot-Product
final_score += sycl::dot(q_float_vec, k_float_vec);
}
return final_score;
} else {
// Fallback für float
float score = 0.0f;
for (int di = 0; di < d_k; ++di) {
score += Q_row_float[di] * K_ptr[di];
}
return score;
}
}
// Haupt-Flash-Attention Kernel
/**
/
template
void flash_attention_kernel_impl(
const scalar_t Q_ptr,
const scalar_t* K_ptr,
const scalar_t* V_ptr,
scalar_t* Out_ptr,
int num_q,
int num_k,
int d_k,
int d_v,
int q_stride,
int k_stride,
int v_stride,
int out_stride,
sycl::nd_item<1> item
) {
const int head_row = item.get_global_id(0);
if (head_row >= num_q) return;
// Lokale Register/Akkumulatoren
float accum_den = 0.0f; // Denominator (Z)
float running_max = -INFINITY; // Globaler Max-Score (m)
// Temporärer Speicher für Scores (in Registern/private memory)
float S_scores[BLOCK_N];
float accum_num[D_MAX] = {0.0f}; // Numerator (P*V Summe)
// Lade gesamte Q-Zeile in Float-Register
const scalar_t* Q_row_ptr = Q_ptr + head_row * q_stride;
float Q_row_float[D_MAX];
for (int di = 0; di < d_k; ++di) {
Q_row_float[di] = static_cast(Q_row_ptr[di]);
}
const float scale_factor = 1.0f / sycl::sqrt(static_cast(d_k));
// Tiling-Strategie: Iteration über K/V-Blöcke
for (int k_start = 0; k_start < num_k; k_start += BLOCK_N) {
const int k_block_size = sycl::min(BLOCK_N, num_k - k_start);
float block_max = running_max;
// --------------------------------------------------------------------
// Phase 1: Scores berechnen und Max-Score finden
// --------------------------------------------------------------------
for (int kk = 0; kk < k_block_size; ++kk) {
const int k_idx = k_start + kk;
const scalar_t* K_block_ptr = K_ptr + k_idx * k_stride;
// Dot-Product berechnen
float score = dot_product_vec(Q_row_float, K_block_ptr, d_k);
score *= scale_factor;
S_scores[kk] = score; // Score-Caching
// Update Maximum
block_max = sycl::fmax(block_max, score);
}
running_max = sycl::fmax(running_max, block_max);
// --------------------------------------------------------------------
// Phase 2: Softmax-Akkumulation
// --------------------------------------------------------------------
float exp_sum = 0.0f;
for (int kk = 0; kk < k_block_size; ++kk) {
float exp_val = sycl::exp(S_scores[kk] - running_max);
// Numerator-Akkumulation
const scalar_t* V_block_ptr = V_ptr + k_start * v_stride;
for (int dv = 0; dv < d_v; ++dv) {
accum_num[dv] += Q_row_float[dv] * V_block_ptr[k_start * d_v + dv];
}
// Denominator-Akkumulation
accum_den += exp_val;
// Update Scores (gemäß gemitteltem Softmax)
S_scores[kk] = exp_val;
}
// --------------------------------------------------------------------
// Phase 3: Output-Akkumulation
// --------------------------------------------------------------------
const float alpha = running_max; // Alpha = log(Z)
for (int kk = 0; kk < k_block_size; ++kk) {
const int k_idx = k_start + kk;
float exp_val = S_scores[kk] * alpha;
// Akkumulation des Outputs
const scalar_t* V_block_ptr = V_ptr + k_idx * v_stride;
for (int dv = 0; dv < d_v; ++dv) {
Out_ptr[head_row * out_stride + dv] += Q_row_float[dv] * exp_val * V_block_ptr[dv];
}
}
// --------------------------------------------------------------------
// Phase 4: Numerator-Normalisierung
// --------------------------------------------------------------------
float denominator = accum_den;
for (int dv = 0; dv < d_v; ++dv) {
Out_ptr[head_row * out_stride + dv] += accum_num[dv] * scale_factor;
}
// Reset Numerator-Akkumulatoren
for (int dv = 0; dv < d_v; ++dv) {
accum_num[dv] = 0.0f;
}
}
}
// Wrapper-Kernel für Flash Attention (Intel Arc GPU optimiert)
template
class FlashAttentionKernel {
private:
int num_q;
int num_k;
int d_k;
int d_v;
int q_stride;
int k_stride;
int v_stride;
int out_stride;
scalar_t* Q_ptr;
scalar_t* K_ptr;
scalar_t* V_ptr;
scalar_t* Out_ptr;
public:
FlashAttentionKernel(
const scalar_t* Q,
const scalar_t* K,
const scalar_t* V,
scalar_t* O,
int num_queries,
int num_keys,
int head_dim,
int head_dim_v,
int q_stride,
int k_stride,
int v_stride,
int out_stride
) : num_q(num_queries), num_k(num_keys), d_k(head_dim), d_v(head_dim_v),
q_stride(q_stride), k_stride(k_stride), v_stride(v_stride),
out_stride(out_stride), Q_ptr(const_cast<scalar_t>(Q)), K_ptr(const_cast<scalar_t>(K)),
V_ptr(const_cast<scalar_t*>(V)), Out_ptr(O) {}
void operator()() const {
// Get SYCL Device und Context
auto q_device = Q_ptr->get_device();
main: context full and context shift is disabled => stopping
common_perf_print: sampling time = 226,83 ms
common_perf_print: samplers time = 108,39 ms / 2048 tokens
common_perf_print: load time = 9042,39 ms
common_perf_print: prompt eval time = 1855,20 ms / 1257 tokens ( 1,48 ms per token, 677,55 tokens per second)
common_perf_print: eval time = 118635,30 ms / 790 runs ( 150,17 ms per token, 6,66 tokens per second)
common_perf_print: total time = 120739,40 ms / 2047 tokens
common_perf_print: unaccounted time = 22,06 ms / 0,0 % (total - sampling - prompt eval - eval) / (total)
common_perf_print: graphs reused = 0
llama_memory_breakdown_print: | memory breakdown [MiB] | total free self model context compute unaccounted |
llama_memory_breakdown_print: | - SYCL0 (Intel(R) Arc(TM) A730M Graphics) | 11597 = 11597 + (10083 = 9547 + 272 + 264) + 17592186034332 |
llama_memory_breakdown_print: | - Host | 1036 = 1024 + 0 + 12 |
✅->AI/KI-ANTWORT-FERTIG-GLÜCKWUNSCH
🔷 🎯GLÜCKWUNSCH✅XAIGPUARC🧠ANTWORT✨ABGESCHLOSSEN📝UNTERXAIGPUARC/bin/llama-cli
🔷 DER🏗BAUVERLAUF📝VON-XAIGPUARC-WIRD HIER GESPEICHERTXAIGPUARC/XAIGPUARC.log
╭─alu@king in ~ took 2m13s
╰─λ Ab Phase ein hat er selbst geschrieben wie man hier selber auch nachlesen kann.
Die Werte, sind garnicht mal so Falsch... und entsprechen dem Grundbild des wirklichen weiteren Teil des Flash Attention Kernels. Ich habe schlicht nicht den ganzen Kernel eingeladen, weil wir hier noch auf etwa 2000 Token beschränkt sind. Dafür, arbeiten diese aber erstaunlich gut für schlichtes Testen der Leistungsfähigkeit ohne jeglichen Optimierungen auf die Qualität der Antworten. Das genutzt Modell ist auch eines der neuen.
Die letzten Wochen und Monate gab es nicht viel zu sehen am KI Markt, das hat sich aber inzwischen geändert. Sozusagen findet gerade ein irgendwie "aufbäumen" statt. Die guten Modelle findet man teilweise schon nach einem Tag nicht mehr einfach so zum Runterladen.
Man sollte sich KIs JETZT Sichern, später wird es sicherlich entwicklungen geben, aber man darf nicht vergessen, das die Daten dann auch von KIs selbst stammen und eben nicht mehr den nackten Zustand der Menschheit darstellen, wie er eben noch vor zwei Jahren war.
Vor über zwei Jahren, habe ich übrigens die Erste KI auf dem Laptop hier ausprobiert....
Edit: Das hier ist die öffentliche Version von XAIGPUARC!!! Man kann das also nachtesten für sich.
Meine Nichtöffentlichen Versionen sind schon leicht weiter Fortgeschritten, aber nicht Gut genug, das ich schon Veröffentlichen möchte. Sie werden einfach etwas Größere Modelle auch auf CPU und geteilter GPU Berechenbar machen. Womit zwar die Geschwindigkeit massiv leiden wird, aber man eben auch mit "wenig" sehr große Modelle laden können soll. Maximale Modellfreiheit sozusagen in Vollarbeit.
Ich habe noch etwas weiter getestet und muss einfach Sagen, der König der Modelle ist MathTutor und man bekommt ihn nicht ohne Anmeldung. Aber er ist einfach der König, da gibt es nicht viel zu Reden.
Das gilt halt leider nur für die A770 mit ihren 16GB
Aber dort ist sie mit ihm nicht zu schlagen in zufriedenstellender Antwort mit 2K Token im Promt und unter einer Minute Anwortzeit. Keine Schnörkel in den Antworten, einfach brachial das Logische Fortgesetzt. Und das Gut für die wichtigste Sache nämlich Logik.
Mathe kann er bestimmt auch Gut, na Gut, das ist ja auch Logik glaube ich, also Mathe, nicht der Tutor. :-)
Die neuen Modelle wissen ziemlich viel und ja, dafür kann man die schon hernehmen, zum Arbeiten in Form von Coden würde ich die aber nur als Prüfer einsetzen.
Ich meine das ist schon irre hier zu reden das ich bald ohne Netz Vibe Coden kann und ein ernsthafter Konkurrent werde im Kernelbau mit Mathezeug. ;-) Ich bau mir mein Werkzeug, um mir Werkzeuge für das Bauen zu bauen. :-D