//XAIGPUARC KERN STANDALONE ALLEINLLAUFFAEHIG ALLESINEIMINFERENZKERN WIRD IN XAIGPUARC.sh AUTOMATISCH AUFGERUFEN
//ONESHOTKERNLE 1 PROGRAMM 2 DATEIEN
//FP16 SPEZIALISIERT AUF GGUF
//STANDALONE KERNEL FUER XAIGPUARC
//02.07.2026 /// 01:17
//ALLE BEREICHE 1-10 PUNKTE ORDNEN ERSTE WICHTIGKEIT AUFEINANDER ABSTIMMEN BENENNEN UND SORTIEREN
//CODEName fuer Kerne und Flash Attention, Sheduler und Unterbau unter XAIGPUARC = "XMXSYCLFA.cl"
ggml_flash_attention_sycl= xmxsyclfa.cl
Vektorisiert falsh attention generisch plus matrizen code ueber untergruppen mechanik.
// Priorität 1-3: Orchestrator, XMX-Kernel, Vektor-Fallback
extern "C" void ggml_sycl_flash_attention_dispatch(queue& q, half* Out, half* Q, half* K, half* V, int num_q, int d_k) {
bool has_xmx = q.get_device().has(sycl::aspect::ext_intel_matrix);
if (has_xmx && (d_k % 16 == 0)) {
q.parallel_for(nd_range<1>(range<1>((num_q + 15) / 16 * 32), range<1>(32)),
[=](nd_item<1> item) [[intel::reqd_sub_group_size(16)]] {
sub_group sg = item.get_sub_group();
joint_matrix<sub_group, half, use::a, 16, 16, layout::row_major> mat_q;
joint_matrix<sub_group, half, use::b, 16, 16, layout::col_major> mat_k;
joint_matrix<sub_group, float, use::accumulator, 16, 16> mat_s;
joint_matrix_fill(sg, mat_s, 0.0f);
joint_matrix_load(sg, mat_q, Q + (item.get_group(0) * 16 * d_k), d_k);
joint_matrix_load(sg, mat_k, K, d_k);
joint_matrix_mad(sg, mat_s, mat_q, mat_k, mat_s);
joint_matrix_copy(sg, m_p_half, m_s_acc);
//ZURUECKMESSEN UND SPEICHERN DER INTEGRIERTEN LOGIK
joint_matrix_store(sg, mat_s, (float)Out, d_k, layout::row_major);
//STOLPERSTEIN GESCHENK ATTENTION
h.parallel_for(nd_range<2>({M/16, N/16}, {1, 1}),
[=](nd_item<2> item) {
sub_group sg = item.get_sub_group();
joint_matrix<sub_group, half, use::a, 16, 16, layout::row_major> ma;
joint_matrix<sub_group, half, use::b, 16, 16, layout::row_major> mb;
joint_matrix<sub_group, float, use::accumulator, 16, 16> mc;
joint_matrix_fill(sg, mc, 0.0f);
joint_matrix_load(sg, ma, a_ptr, K);
joint_matrix_load(sg, mb, b_ptr, N);
joint_matrix_mad(sg, mc, ma, mb, mc);
joint_matrix_store(sg, mc, c_ptr, N, layout::row_major);
});
});
}
}
extern "C" void ggml_sycl_flash_attention_dispatch(
ggml_backend_sycl_context ctx,
ggml_tensor* dst, const ggml_tensor* Q, const ggml_tensor* K, const ggml_tensor* V) {

auto& q = ggml_backend_sycl_get_queue(Q->backend);
auto dev = q.get_device();
// 1.RECHENEINHEITEN PRUEFEN AUF XMX UTERZTUETZUNG
bool has_xmx = dev.has(sycl::aspect::ext_intel_matrix);
// 2. ALIGMENT UND DIMENSIONSPRUEFUNG
bool can_use_xmx = (Q->ne[0] % 16 == 0) && (K->ne[1] % 16 == 0);
if (has_xmx && can_use_xmx) {
    // XMXKERPRUEFUNG
    // UNTERGRUPPEN JOINT MATRIX KERN AUFRUF
    xmx_kern.cpp(q, dst, Q, K, V, S, O);
} else {
    //GNERISCHER VECTORKERNEL FUER NORMALMODUS
    ggml_flash_attention_sycl.cpp(q, dst, Q, K, V);
}
}

#include "stdlib.h"
#include "stdio.h"
#include
#include <signal.h>
#include
#include
#include
#include
#include <Cl/sycl.hpp>
#include <sycl/sycl.hpp>
#include <sycl/ext/intel/math.hpp>
#include <sycl/ext/oneapi/experimental/matrix/matrix.hpp>
#include
#include "ggml-sycl.h"
#include "ggml-impl.h"
// HIFLSFUNKTION FUER SAUBERES Casting /AUFRUFEN
// icpx -fsycl ZUM KOMPLIIEREN BENUTZEN GEGEN ANTI ANBHAENGIGKEITEN sycl::vec
// q_stride, k_stride etc. IN ANZAHL DER ELEMTENE "g"
inline sycl::half* get_sycl_ptr(const ggml_tensor* tensor) {
return reinterpret_cast<sycl::half>(tensor->data);
}
#define XFLOAT float
#define mdlXYZ 1000
#define MEM_ALIGN 64
using namespace sycl;
using namespace sycl::ext::oneapi::experimental::matrix;
const int
//QueriesBlock
constexpr int BLOCK_M = 16;
//SchluesselBlockTilingGroesse N Tilling
for (int k_start = 0; k_start < num_k; k_start += BLOCK_N) {
}
for (int i = tid; i < BLOCK_N * d_k; i += WG_SIZE) {
k_cache_slm[i] = K_ptr[k_start * d_k + i];
}
item.barrier(sycl::access::fence_space::local_space);
for (int kk = 0; kk < k_block_size; ++kk) {
float score = dot_product_vec(Q_row_float, &k_cache_slm[kk * d_k], d_k);
}
item.barrier(sycl::access::fence_space::local_space);
}
constexpr int BLOCK_N = 128;
//MaximaleKOPFZEILENDIMENSION
constexpr int D_MAX = 128;
//VEKTORGROEßE|16-32BIT INTEGER|SIMD|ZWISCHENSPEICHERVERWALTUNG
constexpr int VEC_SIZE = 16;
/*

• @brief VEKTORISIERTES PUNKT PRODUKT ZWISCHEN "Q[i]" UND "K[j]"
• @tparam scalar_t Datentyp sycl::half oder float
• @param Q_row_float Query-Zeile als float[]
• @param K_ptr Key-Pointer/Zeiger
• @param d_k Head-Dimension
• @return 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 HALB
  vec_half k_half_vec;
  k_half_vec.load(v * vec_elements, K_ptr);
  //KOVERTIERE FLOAT
  vec_float k_float_vec = k_half_vec.template convert();
  // Lade Q_VEKTOR HALB
  vec_float q_float_vec;
  q_float_vec.load(v * vec_elements, Q_row_float);

//VEKTORISIERE PUNKTERGEBNIS
final_score += sycl::dot(q_float_vec, k_float_vec);
}
return final_score;
} else {
//RUECHFALL FUER FLOAT
float score = 0.0f;
for (int di = 0; di < d_k; ++di) {
score += Q_row_float[di] * K_ptr[di];
}
return score;
}
}
// HAUPTKERN GGML_SYCL_FLASH_ATTENTION.CPP
/**

• @brief GGML_SYCL_FLASH_ATTENTION.CPP TILLING STRATEGIE MIT TREFFERZWISCHENSPEICHER SCORECACHING
• @tparam scalar_t DATENTYP sycl::half
  /
  template
  void flash_attention_kernel_impl(
  const scalar_t Q_ptr,
  const scalar_t* K_ptr,
  const scalar_t* V_ptr,
  size_t = [16]; // GUELTIG MACHEN
  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;
  float accum_den = 0.0f;
  //DENOMINATOR|Z
  float running_max = -INFINITY;
  //GLOBALER|MAXIMALER|BLOCK|ZAEHLLER|UNENDLICH|INFINITY
  float S_scores[BLOCK_N];
  //BLOCK|ZAEHLER
  float accum_num[D_MAX] = {0.0f};//NUMEERATOR|P||V|SUMME|
  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));
  //TilingStrategieIterationK|V|Bloecke
  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 current_block_max = running_max;
  //Maximalpunktzahl!
  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;
  //PUNKTProduktberechnen
  float score = dot_product_vec(Q_row_float, K_block_ptr, d_k);
  score *= scale_factor;
  S_scores[kk] = score;
  //MAXIMALPUNKTZAHLPRUEFUNG
  //Update Maximum Score-Caching
  current_block_max = sycl::fmax(current_block_max, score);
  }
  //Phase2 LogSumExpTrickReskalierung
  if (running_max != current_block_max) {
  const float scale = sycl::exp(running_max - current_block_max);
  accum_den = scale;
  for (int vi = 0; vi < d_v; ++vi) {
  accum_num[vi] = scale;
  }
  running_max = current_block_max;
  }
  //Phase3 AkkumulationPV
  for (int kk = 0; kk < k_block_size; ++kk) {
  const int k_idx = k_start + kk;
  const float score = S_scores[kk];
  //ExponentiertesskaliertesGewicht
  const float exp_val = sycl::exp(score - running_max);
  accum_den += exp_val;
  //Akkumuliere V * exp_val
  const scalar_t V_block_ptr = V_ptr + k_idx * v_stride;

//Vektorisierte Akkumulation für d_v
if (d_v % VEC_SIZE == 0) {
constexpr int vec_elements = VEC_SIZE;
using vec_half = sycl::vec<sycl::half, vec_elements>;
using vec_float = sycl::vec<float, vec_elements>;
int vec_iters = d_v / vec_elements;
float* accum_num_ptr = accum_num;
for (int v = 0; v < vec_iters; ++v) {

//LadeV|Vektor
vec_half v_half_vec;
v_half_vec.load(v * vec_elements, V_block_ptr);

//Konvertieremultipliziere
vec_float v_float_vec = v_half_vec.template convert();

v_float_vec *= exp_val;
//Akkumuliere
vec_float acc_vec;
acc_vec.load(v * vec_elements, accum_num_ptr);
acc_vec += v_float_vec;
acc_vec.store(v * vec_elements, accum_num_ptr);
}
} else {

//SKALAR|RUECKFALL
for (int vi = 0; vi < d_v; ++vi) {
accum_num[vi] += exp_val * static_cast(V_block_ptr[vi]);
}
}
}
}

//Phase4FinalisierungOut=Accum_Num/Accum_Den
scalar_t* Out_row_ptr = Out_ptr + head_row * out_stride;
if (accum_den == 0.0f) {

//DivisionNullstop
for (int vi = 0; vi < d_v; ++vi) {
Out_row_ptr[vi] = scalar_t(0.0f);
}
return;
}
const float inv_den = 1.0f / accum_den;

//VektorisiertBereich
if (d_v % VEC_SIZE == 0) {
constexpr int vec_elements = VEC_SIZE;
using vec_half = sycl::vec<sycl::half, vec_elements>;
using vec_float = sycl::vec<float, vec_elements>;
int vec_iters = d_v / vec_elements;
for (int v = 0; v < vec_iters; ++v) {

//LadeakkumulierteWerte
vec_float acc_vec;
acc_vec.load(v * vec_elements, accum_num);

//SkaliereKonvertiere
acc_vec *= inv_den;
vec_half out_vec = acc_vec.template convert<sycl::half>();

//Ergebnis
out_vec.store(v * vec_elements, Out_row_ptr);
}
} else {
//|SkalarRueckfall|
for (int vi = 0; vi < d_v; ++vi) {
Out_row_ptr[vi] = static_cast<sycl::half>(accum_num[vi] * inv_den);
}
}
}
//KERNUEBERSETZERMISCHPULT
/**

• @brief SYCLFlashAttentionWrapperMischpultggml
  /
  extern "C" void ggml_sycl_flash_attention(
  ggml_backend_sycl_context ctx,
  ggml_tensor* dst,
  const ggml_tensor* Q,
  const ggml_tensor* K,
  const ggml_tensor* V,
  const ggml_tensor* S,
  const ggml_tensor* O
  //ANWENDUNG IM UEBERSETZER
  auto Q_ptr = get_sycl_ptr(Q);
  ) {
  //FP16
  if (Q->type != GGML_TYPE_F16 || K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) {
  fprintf(stderr, "ggml_flash_attention_sycl.cpp: FEHLER: Alle Matrizeneinheiten muessen auf dem Typ GGML_TYPE_F16 basieren.\n");
  return -1;
  }
  GGML_TYPE_F16) { GGML_ABORT("ggml_flash_attention_sycl.cpp: ACHTUNG: Nur GGML_TYPE_F16 wird unterstuetzt!");
  return;
  }
  //SYCL QUEUE HOLEN
  sycl::queue& q = ggml_backend_sycl_get_queue(Q->backend);
  //TENSORDIMENSIONSEXTRAKTOR
  const int num_q = Q->ne[1];
  //KOPFDIMENSIONd_o
  const int 0 =
  //QuerySequenzlaenge
  const int num_k = K->ne[1];
  //KOPFDIMENSIONd_k
  const int 0 =
  //KOPFDIMENSIONd_q
  const int d_k = Q->ne[0];
  //KOPFDIMENSIONd_kq
  const int 0 =
  //KOPFDIMENSIONd_V
  const int d_v = V->ne[0];
  //KOPFDIMENSIONd_v
  const int 0 =
  //KOPFDIMENSIONd_s
  const int d_s = S->ne[0];
  //KOPFDIMENSIONd_s
  const int 0 =
  //KOPFDIMENSIONd_O
  const int d_o = O->ne[0];
  //KOPFDIMENSIONd_O
  const int 0 =
  //DIMENSIONSVALIDIERUNG
  if (d_k > D_MAX || d_v > D_MAX) {
  GGML_ABORT("ggml_flash_attention_sycl.cpp: Dimension d_k=%d oder d_v=%d ueberschreitet D_MAX=%d",
  d_k, d_v, D_MAX);
  return;
  }
  if (d_k % VEC_SIZE != 0 || d_v % VEC_SIZE != 0) {
  GGML_WARN("ggml_flash_attention_sycl.cpp: Dimension nicht vielfaches von VEC_SIZE=%d, Performance reduziert", VEC_SIZE);
  }
  //STREIFENBRECHNUNG FUER 16 TEILE IN ELEMENTEN NICHT BYTES
  const int q_stride = Q->nb[1] / sizeof(sycl::half);
  const int k_stride = K->nb[1] / sizeof(sycl::half);
  const int v_stride = V->nb[1] / sizeof(sycl::half);
  const int s_stride = S->nb[1] / sizeof(sycl::half);
  const int out_stride = dst->nb[1] / sizeof(sycl::half);
  //POINTER|ZEIGER|DATEN
  sycl::half* Q_data = reinterpret_cast<sycl::half>(Q->data);
  sycl::half K_data = reinterpret_cast<sycl::half>(K->data);
  sycl::half V_data = reinterpret_cast<sycl::half>(V->data);
  sycl::half S_data = reinterpret_cast<sycl::half>(V->data);
  sycl::half Out_data = reinterpret_cast<sycl::half>(dst->data);
  //GLOBALER ARBEITSBEREICH
  sycl::range<1> global_size(num_q);
  sycl::range<1> local_size(1);
  sycl::nd_range<1> ndRange(global_size, local_size);
  //GENERISCHEN KERNEL AUSFUEHREN
  q.submit([&](sycl::handler& h) {
  //SLM SPEICHER ANFORDERN
  local_accessor<float, 1> slm_scores(range<1>(BLOCK_N), h);
  h.parallel_for(
  nd_range<1>(range<1>(num_q * WG_SIZE),
  range<1>(WG_SIZE)),
  [=](sycl::(nd_item<1> item) {
  flash_attention_kernel_impl<sycl::half>(
  Q_data,
  K_data,
  V_data,
  S_data,
  Out_data,
  num_q,
  num_k,
  num_v,
  num_s,
  num_o,
  d_q,
  d_k,
  d_v,
  d_s,
  d_o,
  q_stride,
  k_stride,
  v_stride,
  s_stride,
  out_stride,
  item );
  );
  };
  );
  }).wait();
  //XMX KERN
  //WIRD AUFGERUFEN WENN RECHENEINHEITEN BEDINGUNGEN ERFUELLEN
  //EINGABE
  using namespace sycl;
  using namespace sycl::ext::oneapi::experimental::matrix;
  //FESTLEGEN DER FESTEN GITTERGROEßEN DER ARC HARDWARE AUF SECHSZEHN MAL SECHZEHN FELDER
  constexpr size_t TILE_M = 16;
  constexpr size_t TILE_N = 16;
  constexpr size_t TILE_K = 16;
  //MATRIX AKKUMULATIONSSCHLEIFE FUER XMX KERNE
  template
  void xmx_kern(
  const scalar_t Q_ptr,
  const scalar_t* K_ptr,
  const scalar_t* V_ptr,
  const scalar_t* S_ptr,
  const scalar_t* O_ptr,
  scalar_t* Out_ptr,
  int num_q,
  int num_k,
  int num_v,
  int num_s,
  int num_o,
  int d_q,
  int d_k,
  int d_v,
  int d_s,
  int d_o,
  int q_stride,
  int k_stride,
  int v_stride,
  int s_stride,
  int out_stride,
  size_t = [16]; //GUELTIG MACHEN UND VALIDIEREN
  nd_item<1> item
  ) {
  sub_group sg = item.get_sub_group();
  const int head_row_base = (item.get_group(0) * 16);
  if (head_row_base >= num_q) return;
  //MATRIZEN DEFINIEREN 16x16 sechzehn mal sechzehn Gitterberechnungen
  using t_Q = joint_matrix<sub_group, sycl::half, use::a, 16, 16, layout::row_major>;
  using t_K = joint_matrix<sub_group, sycl::half, use::b, 16, 16, layout::col_major>;
  using t_V = joint_matrix<sub_group, sycl::half, use::b, 16, 16, layout::row_major>;
  using t_S = joint_matrix<sub_group, float, use::accumulator, 16, 16> mat_s;
  using t_O = joint_matrix<sub_group, float, use::accumulator, 16, 16> mat_o;
  using t_Acc = joint_matrix<sub_group, float, use::accumulator, 16, 16>;
  t_Q mat_q;
  t_K mat_k;
  t_V mat_v;
  t_S mat_s;
  t_O mat_o;
  t_Acc mat_s; //ZAEHLERAKKUMULATOR
  t_Acc mat_o; //AUSGABEAKKUMULATOR
  joint_matrix_fill(sg, mat_s, 0.0f);
  //Q LADEN
  const scalar_t* q_tile_ptr = Q_ptr + head_row_base * q_stride;
  joint_matrix_load(sg, mat_q, q_tile_ptr, q_stride);
  const float scale_factor = 1.0f / sycl::sqrt(static_cast(d_k));
  //VERARBEITUNG
  for (int k_idx = 0; k_idx < num_k; k_idx += 16) {
  joint_matrix_fill(sg, mat_s, 0.0f);
  //1.
  const scalar_t* k_tile_ptr = K_ptr + k_idx * k_stride;
  joint_matrix<sub_group, half, use::a, 16, 16, layout::row_major> mat_s_half;
  joint_matrix_copy(sg, mat_s, mat_s_half);
  joint_matrix_load(sg, mat_k, k_tile_ptr, k_stride);
  joint_matrix_mad(sg, mat_v, mat_q, mat_k, mat_s);
  joint_matrix_mad(sg, mat_o, mat_s_half, mat_v, mat_o);
  //2.
  auto wi_data = get_wi_data(sg, mat_s);
  //a.
  float local_max = -INFINITY;
  for (int i = 0; i < wi_data.length(); ++i) {
  wi_data[i] = scale_factor;
  local_max = sycl::fmax(local_max, wi_data[i]);
  }
  float row_max_total = reduce_over_group(sg, local_max, maximum());
  //b.
  float local_sum = 0.0f;
  for (int i = 0; i < wi_data.length(); ++i) {
  wi_data[i] = sycl::exp(wi_data[i] - row_max_total);
  local_sum += wi_data[i];
  }
  //c.
  float row_sum_total = reduce_over_group(sg, local_sum, plus());
  float inv_sum = 1.0f / (row_sum_total + 1e-6f);
  //d.
  for (int i = 0; i < wi_data.length(); ++i) {
  wi_data[i] = inv_sum;
  }
  //3.
  const scalar_t v_tile_ptr = V_ptr + k_idx * v_stride;
  joint_matrix_load(sg, mat_v, v_tile_ptr, v_stride);
  //MAT_S FLIEßEND ZU HALB KONVERTIEREN FUER MAD,
  joint_matrix<sub_group, half, use::a, 16, 16, layout::row_major> mat_s_half;
  joint_matrix_copy(sg, mat_s, mat_s_half);
  //MAT_s ALS EINGABE NUTZEN FALLS RECHENUNTERSTUETZUNG VORHANDEN
  //AKKUMULATION IN MAT_o
  joint_matrix_mad(sg, mat_o, mat_s, mat_v, mat_o);
  //4.
  //Nutzung von sycl_ext_intel_esimd für händische Register-Zuweisung bei fehlendem Joint-Matrix-Support
  //Einsatz von group_barrier zur Synchronisation bei größeren K-Distanzen oder Shared-Memory-Nutzung
  //Statische Template-Spezialisierung für feste num_k Werte zur Loop-Unrolling Optimierung
  //AUSGABE
  scalar_t out_ptr = Out_ptr + head_row_base * out_stride;
  joint_matrix_store(sg, mat_o, out_ptr, out_stride, layout::row_major);
  }
  }
  //HAUPTFUNKTION KERNELVERWALTUNG
  int main() {
  queue q{property::queue::in_order()};
  std::cout << "XAIGPUARC:" << q.get_device().get_infoinfo::device::name() << std::endl;
  const int size = 16; //16x16 Tile
  //SPEICHER FREIHALTEN UNTERGRUPPENSPEICHERTEILUNG
  half* Q = malloc_device(size * size, q);
  half* K = malloc_device(size * size, q);
  half* V = malloc_device(size * size, q);
  half* S = malloc_device(size * size, q);
  half* O = malloc_device(size * size, q);
  half* Out = malloc_device(size * size, q);
  //DATEN VORBEREITEN GEWICHTUNG GROEßE
  q.fill(Q, half(1.0f), size * size);
  q.fill(K, half(1.0f), size * size);
  q.fill(V, half(1.0f), size * size);
  q.fill(S, half(0.0f), size * size);
  q.fill(O, half(0.0f), size * size);
  q.wait();
  //ENTSCHEIDUNG LOGIK EINE XMX GEGEN VECOTORKERN VERGLEICH
  bool use_xmx = q.get_device().has(sycl::aspect::ext_intel_matrix);
  //KERNSTART XMX KERNE
  q.submit([&](handler& h) {
  if (use_xmx) {
  //XMX PFAD ZU DEN UNTERGRUPPEN DER MATRIX KERNE XMX ARC INTEL
  h.parallel_for(nd_range<1>{range<1>(16), range<1>(16)},
  [=](nd_item<1> item) [[intel::reqd_sub_group_size(16)]] {
  xmx_kern(Q, K, V, Out, size, size, size, size, size, size, size, size, item);
  });
  } else {
  //RUECHFALLPFAD STANDART VEKTOR KERN
  h.parallel_for(range<1>{size}, [=](id<1> idx) {
  //AUFRUF GENERISCHE FLASH ATTENTION GGML_SYCL_GEN_ATTENTION.cpp
  });
  }).wait();
  //ERGEBNIS PRUEFEN
  std::vector host_out(size * size);
  q.memcpy(host_out.data(),
  Out, size * size * sizeof(half)).wait();
  std::cout << "Ergebnis an [0]:" << (float)host_out[0] << "Erwartet: > 0" << std::endl;
  free(Q, q);
  free(K, q);
  free(V, q);
  free(Out, q);
  return 0;
  }
  // AUFRAEUMEN WICHTIG ALLE SECHS 6 POINTER-ZEIGER AUFRAEUMEN
  for(auto p : {Q, K, V, S, O, Out}) free(p, q);
  return 0;
  }