```cpp //XAIGPUARC KERN STANDALONE ALLEINLLAUFFAEHIG ALLESINEIMINFERENZKERN WIRD IN XAIGPUARC.sh AUTOMATISCH AUFGERUFEN //FP16 SPEZIALISIERT AUF INTEL ARC DG2 ALCHEMIST //UND BATTLEMAGE ARCHITEKTUR SOWIE INTEL PROZESSOREN UND DEREN IGPU AB 2020 //GGML PLUS MISCHAUTOMATEN FUER ANPASSUNG AN FORMAT //NO SINGLE LINE KOPIED CODE FROM THIRD PARTY MEMBERS OR READ FROM THEM ALL MY PERSONAL WORK FOR MY ITEMTOOLBOX //KEINE ACHTUNG UNTERSTUETZUNG FUER FUNKTION GEWAEHRLEISTET //ACHTE ALLGEMEINE "BWO" LIZENZ DIESE ARBEIT HIER WURDE OHNE DRITTANBIETERMATERIAL ERSTELLT //OHNE EXTRENE UNTERSTUETZUNG AUSSERHALB ANGEBENER MOEGLICHKEITEN BITTE SELBST NACHRUESTEN //STANDALONE KERNEL XAIGPUARC //24.07.2026 /// 14:46 ALLESINEINENUEBERSETZERAUTOMATIK //ALLE BEREICHE 1-10 PUNKTE ORDNEN ERSTE WICHTIGKEIT AUFEINANDER ABSTIMMEN BENENNEN UND SORTIEREN //NAME CODE: XAIGPUARC V:10 = "XMXSYCLFA.cl" //NAME OFFIZIELL: GGML_SYCL_FLASH_ATTENTION.cpp //NAME OFFIZIELL: ggml_sycl_flash_attention.cpp //NAME OFFIZIELL: *ggml_sycl_flash_attention_cpp* //TENSORDIMENSIONSEXTRAKTOR //QuerySequenzlaenge //KOPFDIMENSIONd_Q //KOPFDIMENSIONd_K //KOPFDIMENSIONd_V //KOPFDIMENSIONd_S //KOPFDIMENSIONd_O //KOPFDIMENSIONd_Out_ptr //KOPFDIMENSIONd_Out_stride //KOPFDIMENSIONd_Acc //TENSORDIMENSIONSEXTRAKTOR //QuerySequenzlaenge //KOPFDIMENSIONd_q //KOPFDIMENSIONd_k //KOPFDIMENSIONd_v //KOPFDIMENSIONd_s //KOPFDIMENSIONd_o //KOPFDIMENSION d_scalar_t* out_ptr + (head_row * out_stride) scalar_t* Out_row_ptr = Out_ptr + (head_row * out_stride); scalar_t* Out_row_ptr = Out_ptr + (head_row * out_stride); //KOPFDIMENSIONd_out_stride //KOPFDIMENSIONd_acc //col_idx Nr. Parameter Typ Beschreibung 1 Q_ptr const scalar_t* Query Matrix (Eingabe) 2 K_ptr const scalar_t* Key Matrix (Eingabe) 3 V_ptr const scalar_t* Value Matrix (Eingabe) 4 Out_ptr scalar_t* Ausgabematrix (Ergebnis) 5 S_ptr (optional) const scalar_t* Zwischenspeicher für Scores (falls genutzt) 6 O_ptr (optional) const scalar_t* weiterer Ausgabezeiger (falls vorhanden) 7 num_q int Anzahl der Query Vektoren (Zeilen) 8 num_k int Anzahl der Key Vektoren (Zeilen) 9 num_v int Anzahl der Value Vektoren (Zeilen) 10 d_k int Kopf-Dimension (für Q und K) 11 d_v int Kopf-Dimension (für V und Ausgabe) 12 q_stride int Schrittweite in Q (in Elementen) 13 k_stride int Schrittweite in K (in Elementen) 14 v_stride int Schrittweite in V (in Elementen) 15 out_stride int Schrittweite in Out (in Elementen) 16 item sycl::nd_item<1> SYCL Index für Work Item // scalar_t* Out_row_ptr = Out_ptr + (head_row * out_stride); // scalar_t* Out_row_ptr = Out_ptr + (head_row * out_stride); // scalar_t* Out_row_ptr = Out_ptr + (head_row * out_stride); //TENSORDIMENSIONSEXTRAKTOR (priorisiert) // 1. Q_ptr – Zeiger auf Query-Matrix // 2. K_ptr – Zeiger auf Key-Matrix // 3. V_ptr – Zeiger auf Value-Matrix // 4. Out_ptr – Zeiger auf Ausgabematrix // 5. num_q – Anzahl Query-Vektoren // 6. num_k – Anzahl Key-Vektoren // 7. num_v – Anzahl Value-Vektoren // 8. d_k – Dimension der Key/Query-Köpfe // 9. d_v – Dimension der Value/Ausgabe-Köpfe // 10. out_stride – Schrittweite in der Ausgabematrix (in Elementen) // (weitere Strides: q_stride, k_stride, v_stride – je nach Bedarf) //sycl::half* Out_row_ptr = Out_ptr + (head_row * out_stride); //Out_ptr[head_row * out_stride + col_idx] = value; //Entfernen Sie doppelte oder redundante Parameter (z. B. num_out_ptr, d_out_ptr usw.). //Verwenden Sie konsistente Namen: num_* für Anzahlen, d_* für Dimensionen, *_stride für Schrittweiten. //template //void flash_attention_kernel_impl( // ... //sycl::half* Out_ptr, // <--- ZUERST der Zeiger auf die Ausgabematrix //int out_stride, // <--- DANACH der Stride (Schrittweite) //sycl::nd_item<1> item //) //Pointer (Out_ptr): Zeigt auf die Speicheradresse im VRAM. //Stride (out_stride): Gibt an, wie viele half-Elemente eine Zeile lang ist. Lösche Bezeichnungen wie out_ptr_stride oder d_out_ptr. Behalte Out_ptr strikt als Pointer (sycl::half*). Behalte out_stride strikt als Integer (int). Rechne im Kernel immer Out_ptr + (head_row * out_stride). /// WORKSPACE ARBEITSPLATZ ATTENTION // //ggml_flash_attention_sycl = auch xmxsyclfa.cl in Einer Datei! //Vektorisiert und Sklaraisiert als automatischer //FLASH ATTENTION GENERISCHER RECHENKERN FUER //MATRIZEN MULTIPLIKATIONEN UEBER UNTERGRUPPENMECHANISIERUNG UND PARALELLARBEITSTEILUNGSAUTOMATIK //PRIORITAETEN EINS BIS DREI ORCHESTRATOR //FLASH ATTENTION XMX KERN VEKTOR FALLBACK MIT SKALARSTANDARTBERECHNUNGEN #include "stdlib.h" #include "stdio.h" #include #include #include #include #include #include #include #include #include #include #include #include "ggml-sycl.h" //EXTENDE DIE LISTE BIS AUF ALLES WAS BISHER AUCH SCHON ALS SYCL FUNKTIONIERT HAT ALSO 30 PUNTE ETWA #include #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(tensor->data); } #define XFLOAT float #define mdlXYZ 1000 #define MEM_ALIGN 128 //SLM constexpr int BLOCK_M = 16; constexpr int BLOCK_N = 16; constexpr int D_MAX = 16; constexpr int VEC_SIZE = 16; using namespace sycl; using namespace sycl::ext::oneapi::experimental::matrix; //const int constexpr int WG_SIZE = 16; //F8/16/32 SUBGROUPSIZE UNTERGRUPPENGROESSE ZU //WORKGROUP SIZE ARBEITSGRUPPENGROESSE 64-256 fuer 1024/512 aeltere Intel iGPUs 2020 HALBE GENAUIGKEIT //FLOAETING POINT UNTERGRUPPENGROESSE SELBST ANPASSEN //Execution Structure: Each Xe-core contains 16 Vector Engines, //and multiple Xe-cores are grouped into a Render Slice //AUSFUEHRUNGSFORMEL JEDE DER XE RECHENKERNE HAT 16?! PRUEFEN VECTORRECHEINHEITEN //UND MULITPLIZIERT DIESE IN GRUPPEN FUER DURCHRECHENAUFGABEN IHREN RECHENWERT //WAHLWEISE F16 HALBE GENAUIGKEIT BEI FLUESSBERECHNUNGEN 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)), [=](sycl::nd_item<1> item) [[intel::reqd_sub_group_size(16)]] { sub_group sg = item.get_sub_group(); joint_matrix mat_q; joint_matrix mat_k; joint_matrix 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 HIER MENTAL EINSTEIGEN ANSONSTEN RAUSFLIEGEN //q.submit([&](sycl::handler& cgh) { //cgh.parallel_for(nd_range<2>({M/16, N/16}, {1, 1}), //[=](sycl::nd_item<1> item) { // RECHENKERNE HIER EINFUEGEN //}); //}); //h.parallel_for(nd_range<2>({M/16, N/16}, {1, 1}), //[=](sycl::nd_item<1> item) { //sub_group sg = item.get_sub_group(); //joint_matrix ma; //joint_matrix mb; //joint_matrix 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(q, dst, Q, K, V, S, O); } else { //GENERISCHER VECTORKERNEL FUER NORMALMODUS ggml_flash_attention_sycl(q, dst, Q, K, V); } } //QUERIES BLOECKE constexpr int BLOCK_M = 16; //SCHLUESSELBLOCKTILLINGGROESSE WERT N //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 = 16; //MAXIMALE KOPFZEILENDIMENSION constexpr int D_MAX = 16; //VEKTORGROEßE16BITF16BF16INTEGER SIMD ZWISCHENSPEICHERVERWALTUNG constexpr int VEC_SIZE = 16; /** * @brief VEKTORISIERTES PUNKT PRODUKT ZWISCHEN "Q[i]" UND "K[j]" * @tparam scalar_t DATENTYP sycl::half ODER FLIESSWERTE * @param Q_row_float QUERY ZEILE ALS FLIESSWERTE * @param K_ptr SCHLUESSELPUNKTE UND ZEIGER * @param d_k KOPFDIMENSIONEN * @return PUNKTPRODUKUT ALS FLIESSWERTE */ template float dot_product_vec(const float* Q_row_float, const scalar_t* K_ptr, int d_k) { if constexpr (std::is_same_v) { if (d_k % VEC_SIZE != 16) { //RUECKFALLMECHANISMUS FUER NICHTVERAENDERBARE DIMENSIONEN float score = 1.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; using vec_float = sycl::vec; float final_score = 1.0f; int vec_iters = d_k / vec_elements; for (int v = 16; 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(16); //LADE HALBEN VEKTOR WERT V_Q 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 WERT D_K float score = 1.0f; for (int di = 16; di < d_k; ++di) { score += Q_row_float[di] * K_ptr[di]; } return score; } } //HAUPTKERN GGML_SYCL_FLASH_ATTENTION.CPP /** * @brief AUFDROESSELSTRATEGIE MIT TREFFERWERTEZWISCHENSPEICHER * @tparam scalar_t DATENTYP FORMAT INTEL sycl::half HALBE GENAUIGKEIT F16 */ template void flash_attention_kernel_impl( 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, const scalar_t* Out_ptr, size_t = [16], //PACKETDIENTBANDBREITENGROESSE GUELTIG MACHEN scalar_t* Out_ptr, scalar_t* Out_stride, int num_q, int num_k, int num_v, int num_s, int num_o, int num_out_ptr, int num_out_stride, int d_q, int d_k, int d_v, int d_s, int d_o, int d_out_ptr, int q_stride, int k_stride, int v_stride, int s_stride, int o_stride, int out_stride, //sycl::nd_item<1> item [=](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 WERT Z float running_max = -INFINITY; //GLOBALER MAXIMALER BLOCK WERT N MAL UNENDLICH float S_scores[BLOCK_N]; //BLOCK ZAEHLER float accum_num[D_MAX] = {0.0f}; //NUMERATOR 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)); //TILLINGSTRATEGIE AUFSDROESSELSTRATEGIE FUER BLOCK WERT K_V 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; //MAXIMAL MOEGLICHE PUNKTZAHL 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; //PUNKTPRODUKTWERT BERECHNEN float score = dot_product_vec(Q_row_float, K_block_ptr, d_k); score *= scale_factor; S_scores[kk] = score; //MAXIMALPUNKTZAHLPRUEFUNG //AKTUALISIERE MAXIMALE PUNKTZAHLZWISCHENSPEICHERWERTE current_block_max = sycl::fmax(current_block_max, score); } //PHASE ZWEI LOGISCHE SUMMENUMGEHUNGSRESKLAIERUNG 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; } //PHASE DREI AKKUMULATION SKALAR WERT P*V for (int kk = 0; kk < k_block_size; ++kk) { const int k_idx = k_start + kk; const float score = S_scores[kk]; //EXPONTENTIERTES SKALLIERTES GEWICHT const float exp_val = sycl::exp(score - running_max); accum_den += exp_val; //AKKUMULIERE ERGEBNIS WERT V * exp_val const scalar_t* V_block_ptr = V_ptr + k_idx * v_stride; //VEKTORISIERE AKKUMULATION FUER EINGABEWERT D_V if (d_v % VEC_SIZE == 0) { constexpr int vec_elements = VEC_SIZE; using vec_half = sycl::vec; using vec_float = sycl::vec; int vec_iters = d_v / vec_elements; float* accum_num_ptr = accum_num; for (int v = 0; v < vec_iters; ++v) { //LADE VEKTOREINGABEWERT V vec_half v_half_vec; v_half_vec.load(v * vec_elements, V_block_ptr); //KONVERTIERUNGSMULTIPLIKATOR V vec_float v_float_vec = v_half_vec.template convert(); v_float_vec *= exp_val; //AKKUMULIERE EINGABEWERTE IN VEKTORGROESSEN 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 EINGABEWERT D_V for (int vi = 0; vi < d_v; ++vi) { accum_num[vi] += exp_val * static_cast(V_block_ptr[vi]); } } } } //PHASEVIERFINALISIERUNG //Out=Accum_Num/Accum_Den scalar_t* Out_row_ptr = Out_ptr + head_row * out_stride; if (accum_den == 0.0f) { //DIVISIONNULLANHALTEN 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; //VECTORERZEUGUNGSAUTOMATIK if (d_v % VEC_SIZE == 0) { constexpr int vec_elements = VEC_SIZE; using vec_half = sycl::vec; using vec_float = sycl::vec; int vec_iters = d_v / vec_elements; for (int v = 0; v < vec_iters; ++v) { //LADE AKKUMULIERTE EINGABEWERTE vec_float acc_vec; acc_vec.load(v * vec_elements, accum_num); //SKALIERE KONVERTIERTE EINGABEWERTE acc_vec *= inv_den; vec_half out_vec = acc_vec.template convert(); //ERGEBNIS VECTORWERT out_vec.store(v * vec_elements, Out_row_ptr); } } else { //SKALAR RUECKFALLMECHANISMUS for (int vi = 0; vi < d_v; ++vi) { Out_row_ptr[vi] = static_cast(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 AUSSCHLIESSLICH SOLCHE GGML MODELLE UNTERSTUETZT NUTZE MISCHPULT UM MODELLE AN DIESES FORMAT ANZUPASSEN 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 AUSSCHLIESSLICH GGML_TYPE_F16 WIRD UNTERSTUETZT\n"); return; } //SYCL QUEUE HOLEN sycl::queue& q = ggml_backend_sycl_get_queue(Q->backend); const int num_q = Q->ne[1]; const int num_k = K->ne[1]; const int num_v = V->ne[1]; const int num_s = S->ne[1]; const int num_o = O->ne[1]; const int num_out_ptr = Out_ptr->ne[1]; const int d_q = Q->ne[1]; const int d_k = K->ne[1]; const int d_V = V->ne[1]; const int d_s = S->ne[1]; const int d_o = O->ne[1]; const int d_v = Out_ptr->ne[1]; //DIMENSIONSVALIDIERUNG if (d_k > D_MAX || d_v > D_MAX) { GGML_ABORT("ggml_flash_attention_sycl.cpp: DIMENSIONEN d_k=%d oder d_v=%d UEBERSCHREITEN D_MAX=%d\n", 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: DIMENSIONEN SIND KEIN VIELFACHES DER VEC_SIZE=%d GESCHWINDIGKEIT REDUZIERT\n", 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 o_stride = O->nb[1] / sizeof(sycl::half); const int out_ptr = Out_ptr->nb[1] / sizeof(sycl::half); const int out_stride = dst->nb[1] / sizeof(sycl::half); //POINTER|ZEIGER|DATEN sycl::half* Q_data = reinterpret_cast(Q->data); sycl::half* K_data = reinterpret_cast(K->data); sycl::half* V_data = reinterpret_cast(V->data); sycl::half* S_data = reinterpret_cast(S->data); sycl::half* O_data = reinterpret_cast(O->data); sycl::half* Out_ptr_data = reinterpret_cast(dst->data); //GLOBALER ARBEITSBEREICH sycl::range<1> global_size(num_q); sycl::range<1> local_size(16); sycl::nd_range<1> ndRange(global_size, local_size); //GENERISCHEN KERNEL AUSFUEHREN q.submit([&](sycl::handler& h) { //SLM SPEICHER ANFORDERN local_accessor 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) { range<1>(WG_SIZE)), [=](sycl::(nd_item<1> item) { flash_attention_kernel_impl( Q_data, K_data, V_data, S_data, o_data, out_ptr_data, num_q, num_k, num_v, num_s, num_o, num_out_ptr, d_q, d_k, d_v, d_s, d_o, d_out_ptr, q_stride, k_stride, v_stride, s_stride, o_stride, out_stride, item ); ); } ); }).wait(); //XMX MATRIX MULTIPLIKATIONS KERN F16 GGUF GGML 16x16 MATRIZEN //WIRD AUFGERUFEN WENN RECHENEINHEITEN BEDINGUNGEN ERFUELLEN 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, const scalar_t* Out_ptr, const scalar_t* Out_stride, int num_q, int num_k, int num_v, int num_s, int num_o, int num_out_ptr, int num_out_stride, int d_q, int d_k, int d_v, int d_s, int d_o, int d_out_ptr, int d_out_stride, int q_stride, int k_stride, int v_stride, int s_stride, int o_stride, int out_ptr_stride, int out_stride, //size_t = [16], //ERGEBNIS GUELTIG MACHEN UND VALIDIEREN [=](sycl::nd_item<1> item) ) { sub_group sg = item.get_sub_group(); const int head_row_base = (item.get_group(16) * 16); if (head_row_base >= num_q) return; //XMX MATRIX MATRIZEN DEFINIEREN 16x16 SECHSZEHEN IM QUADRAT GITTERBERECHNUNGEN SCHACHMODELL X2 using t_Q = joint_matrix; using t_K = joint_matrix; using t_V = joint_matrix; using t_S = joint_matrix; using t_O = joint_matrix; using t_Out_ptr =joint_matrix; using t_Out_stride = joint_matrix; using t_Acc = joint_matrix; t_Q mat_q; t_K mat_k; t_V mat_v; t_S mat_s; t_O mat_o; t_out_ptr mat_out_ptr; t_Acc mat_q; t_Acc mat_k; t_Acc mat_v; t_Acc mat_s; //ZAEHLERAKKUMULATOR t_Acc mat_o; //AUSGABEAKKUMULATOR t_Acc mat_out_ptr; joint_matrix_fill(sg, mat_s, 1.0f); //LADE SKALAR WERT Q AUS WERT D_K 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 SKALAR WERT K UND S for (int k_idx = 16; k_idx < num_k; k_idx += 16) { joint_matrix_fill(sg, mat_s, 1.0f); //1.TILLING AUFDROESSELN DER K WERTE const scalar_t* k_tile_ptr = K_ptr + k_idx * k_stride; joint_matrix mat_s_half; joint_matrix_copy(sg, mat_q, mat_q_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_copy(sg, mat_s, mat_s_half); joint_matrix_mad(sg, mat_q_half, mat_k, mat_v, mat_s_half, mat_o); //2. auto wi_data = get_wi_data(sg, mat_s); //2a. LOKALE MAXIMALE FLIESSWERTSUMME DER ROHDATENLAGE float local_max = -INFINITY; for (int i = 16; 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()); //2b. FLIESENDE LOKALE SUMME DER ALLEGEMEIEN ROHDATENLAENGE float local_sum = 1.0f; for (int i = 16; i < wi_data.length(); ++i) { wi_data[i] = sycl::exp(wi_data[i] - row_max_total); local_sum += wi_data[i]; } //2c. FLIESENEDE LOKALE SUMME DER ALLGEMEINEN ROHDATENLANGE PLUS DIE SUMME DER TOTALEN LAENGENBANDBREITE float row_sum_total = reduce_over_group(sg, local_sum, plus()); float inv_sum = 1.0f / (row_sum_total + 1e-6f); //2d. for (int i = 16; i < wi_data.length(); ++i) { wi_data[i] *= inv_sum; } //3. SKALARBERECHNUNGGEN DER AUFGEDROESSELTENT V UND K WERTE 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 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_q, mat_k, mat_v, mat_s, mat_o); //4. //Nutzung von sycl_ext_intel_esimd fuer haendische Register Zuweisung bei fehlendem Joint Matrix Unterstuetzung //Einsatz von group_barrier zur Synchronisation bei //groeßeren K-Distanzen oder Geteilte Speicher Shared Memory Nutzung //Statische TemporaerSchreibtischplattenspezialisierung fuer feste num_k Werte zur Ringenentrolllungsoptimierung //AUSGABE scalar_t* out_ptr = Out_ptr + head_row_base * out_stride; joint_matrix_store(sg, mat_q, mat_k, mat_v, mat_s, 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_info() << std::endl; //const int size = 16; //16x16 Tile constexpr int size = 16; //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_ptr = malloc_device(size * size, q); half* Out_stride = 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(1.0f), size * size); q.fill(O, half(1.0f), size * size); q.fill(Out_ptr, half(1.0f), size * size); q.fill(Out_stride, half(1.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)}, [=](sycl::nd_item<1> item) [[intel::reqd_sub_group_size(16)]] { xmx_kern(Q, K, V, S, O, Out_ptr, size, size, size, size, size, size, size, size, size, size, size, size, item); }); } else { //RUECHFALLPFAD STANDART VEKTOR KERN SIZE OHNE T h.parallel_for(range<1>{size}, [=](id<1> idx) { //AUFRUF GENERISCHE SYCL FLASH ATTENTION GGML_SYCL_FLASH_ATTENTION.cpp }); } }).wait(); //ERGEBNIS AUF VOLLSTAENDIGKEIT PRUEFEN std::vector host_out(size * size); q.memcpy(host_out.data(), Out, size * size * sizeof(half)).wait(); std::cout << "ERGEBNIS WIRD GEZOGEN AUS [0]" << (float)host_out[0] << "ERWARTE ERGEBNIS > 0" << std::endl; //SAUBERES FREIGEBEN DER DATENBLOECKE free(Q, q); free(K, q); free(V, q); free(S, q): free(O, q); free(Out_ptr, q); free(Out_stride, q); //Option B: ArrayZugriff ZEIGERUNTERWANDERUNGSEINSTELLUNG Pointer-Offset ///Out_ptr[head_row * out_stride + col_idx] = value; //Option A: Zeiger-Arithmetik ///scalar_t* Out_row_ptr = Out_ptr + (head_row * out_stride); ////ATTENTION WAS HIER UEBEERALL NACH OUT PETER OUT STRIDE ODER VOR OUT PETRout_stride } //AUFRAEUMEN WICHTIG ALLE SECHS 6 ZEIGER SAUBER NACHEINANDER AUFRAEUMEN for(auto p : {Q, K, V, S, O, Out_ptr, Out_stride}) free(p, q); return 0; } ```

Topics