XAIGPUARC /// Entsyclte Sandkasten KI Version und die Sache mit den Uebersetzungen.
Published on HivePostify by @alucian · Mon Sep 07 2026
Ich uebersetze und uebersetze, die allermeiste Zeit, immer noch Begriffe, die immer wieder zu ein paar wenigen Grundbegriffen zusammenzulaufen scheinen.
Die Paralellen zu meiner Arbeit am Mining Programm, sind immer tiefer fuer mich erkennbar und koennen dadurch auch immer einfacher fuer mich Uebersetzt werden. Natuerlich grundsaetzlich Falsch. Damit ich die Richtige Uebersetzung durch das Ausprobieren aller Fehler entdecke sozusagen. Also sie von Selbst geprueft ist, durch die ganzen Falschen Wege die ich gerne auch mal mit Absicht gehe.
So langsam bekomme ich dadurch den Ueberblick, mir hier wiederum auch Schablonen zu bauen, um die Begriffe zu Normieren und zu Einen, welche ich fuer das Schreiben benoetige.
Rueckwaehrts gewantes Lernen durch das erreichen wollen bestimmter Zielsetzungen, ist dabei auch Gefaehrlich zerstreuend.
Das bedeutet, ich versuche inzwischen die groben Ankerpunkte in den Rechenarbeiten abzustecken und an markanten Stellen bereits entsprechende Begrifflichkeiten immer Gleich zum gemeintem Inhalt zu verweden.
Aber das ist nicht sehr Einfach!!!
Und es dauert sehr Lange, bis ich Fortschritte sehe, vor allem jetzt noch am Anfang, weil die Masse fast Ueberwaeltigend ist.
Aber nur Fast. Denn ich sehe am Ende, immer meinen Computer, den im Kopf meine ich und stecke das entsprechend dort ab, was ich machen will, der ist mein Anker, um mich nicht zu Verlieren in der Masse der Informationen, wie ich ihn Steuern kann und warum ich ihn genauso und nicht anders steuern sollte.
Die Syntax, Sprache, der Menschen in Recheneinheiten, ist dabei das Schwerste fuer mich zu Lernen. Weniger, wie etwas Funktioniert, vielmehr, wie es heisst und warum so.
Wie etwas Funktioniert, kann ich mir Gut vorstellen und Nachvollziehen, wenn ich es als Schaltung betrachte.
Die Worte und Befehle von uns Menschen also am Ende als reine Schaltung zu betrachten und Logikmusterketten fuer entsprechende Funktionen zu bauen, wie im Schlaf, ist eines der Nebenprodukte, die ich Anstrebe, mit meiner chaotischen Art zu Schaffen.
Dabei, erfinde, oder besser entdecke ich dann Funktionen, die ich mir schlicht als Phanasie einfallen lasse.
Auch hier ist die Fehlerkorrektur, bei Scheitern neue Parameter des Weges fuer die Zielerreichung Programm.
Was die Kernprogrammierung der Abfrageautomatiken fuer die Sprachmodelle angeht, habe ich das SYCL Chaos wahrscheinlich so langsam wieder in den Griff bekommen, und hoffe mich Selbst damit zu ueberraschen, das ein oder zwei dieser Fehler, zur Funktion wurden.
Im Grunde mache ich das wie ein Sprachmodell...
Ich rate das Richtige Ergebnis.
Das kennt man schon von meinen Mining Dingen. :-)
Salve
Alucian CPP //XACsyclspielewieseki.cpp //ALLEINSTELLUNGSFUNKTION //XAIGPUARC OpenSource MIT LIZENZ //FROM HERZ WITH HELP OF AI ARCHITEKT alucian Berlin-Buch 07.09.2026 /// 07::55 //Nur F16/(Nebengewicht F32) Ausschliesslich empfohlene Modell in UNTERSTUETZUNG 3 in 1 Kern //SPEZIALISIERT FUER ARC Intel XE++ iGPU+dGPU TENSOR SPLIT ROW XMX NUTZUNG //scalart outrowptr = outptr + (headrow outstride); //XAC ist PRAEFIX fuer XAIGPUARC 128 ASCIII beachten ohne Umlaute und scharfem S //Grossschreibung bevorzugt Zeichenhaltung einfach Nummerierung MAXIMUM 10 Punkte //RICHTIG FALSCH RICHTIG RUECKFALLMECHANISMUS EINBAUEN FUER jointmatrixmad(sg, mats, matq, matk, mats); //CODENAME: SPIELEWIESE KI aka " SANDKASTEN ATTENTION " ERSTLINGSWERK KI KERNBAU NACH OPENCL ARC GPU ZHASH MINING PROGRAMMBAU //XMX SKALAR Softmax (Vor!implementiert) //jointmatrixmad (Q×K^T) (implementiert) //jointmatrixmad (S×V) (implementiert) //scalar RUECKFALLKERNE (implementiert) //XACsycldotproductvec (vektorisiert implementiert) //V-Multiplikation (implementiert) //TEST ATTENTION // ACHTEN IN VERSETZTER FORM ZUERST DIE MARKIERUNG DER 10 ZEHN HAUPTFUNKTIONSABSCHNITTE DANN RECHTS DIE SEKUNDAERABSCHNITTE 0.HILFSFUNKTIONEN // SPIELEWIESE KI RUECKFALLRUECKFALLAUTOMATISIMUSGENERALRUNDIERUNGSACHTUNGSMATRIZENMODUS KLEINSTAUFGABENSPACHMODELL //KernIdentifikation: //Geheimspielsplatzsandkastensprachmodellverortungsidentifikation //IndexAnpassung: //Stellt die BAUGRUNDLAGE FUER XAIGPUARC ALS STANDALONE ALLEINSTELLUNGSKERN DAR UND ERGIBT IN SUMME //ALLE GROSSEN KERNE IN MINIMALFUNKTION SPEZIELL AUF EIN MODELL OPTIMIERT FUER //WENIGE ABER MASSENHAFT UND SCHNELLE AUFGABEN WIE EINFACHE RECHNUNGEN ODER UEBERSETZTUNGEN ZU //ERSTELLEN SOWIE CODESCHNIPSEL ZU TESTEN UND ZU BAUEN //DIESER KERN BEZIEHUNGSWEISE DIE MASSE SEINER UNTER UND NACH DEM XAIGPUARCDateiIndex GERICHTETEN //UNTERKERNEN FUNKTIONIERT AUS SICH SELBST HERAUS UNTER DEN GETESTETEN BEDINGUNGEN BEI ERFOLG IMMER //UND GLEICH IM ERGEBNIS. SEINE FUNKTIONEN SIND ENORM EINGESCHRAENKT ABER IM GEBRAUCH FUER //EINEN GEWISSEN ZWECK HOCHGRADIG OPTIMIERT UND GETESTET. //DIESER BAU STELL LEDIGLICHE EINE FERTIGSTELLUNG VON BEREITS ERREICHTEN UND VALDIERTEN ERGEBNISSEN //MIT EINEM EHEMALIGEM FORK VON LLAMA CPP DAR UND WIRD BIS ZU SEINER VOLLENDUNG GEHEGT UND GEPFLEGT. //DER VORGANG ENTSPRICHT ZU ERSTEINMAL DIESE FOLGENDEN ANSTENDEN PUNKTE AB HIER 1 bis 10 ORDENTLICH //IN DIE RICHTIGE ZEILE EINUFUEGEN!!! PRIORI EINS!!! 1. XACsyclapiattentionqkv //KernIdentifikation: //Aeußeres SYCLZugangsfenster für die QKVKommunikation. //IndexAnpassung: //Entspricht der SchnittstellenDefinition im Bereich kernels/ und bindet sauber an die DispatchEbene an. 2. XACsycldotproductvec //KernIdentifikation: //Vektoriertes Skalarprodukt (F16/F32 Optimierung). //IndexAnpassung: //Uebergang in die VektorKernbibliothek (XACsyclflashattentionvec.h/.cpp). 3. XACsyclflashattentionvecscl //KernIdentifikation: //Skalarer FallbackPunktproduktautomat fuer Flash Attention. //IndexAnpassung: //Abgleich mit dem Pfad kernels/XACsyclflashattentionscl.cpp. 4. XACsyclxmxmatrixopsimpl //KernIdentifikation: //XMXHardwareMatrixoperationen (SubGroups / Intel Arc). //IndexAnpassung: //Ausrichtung an kernels/XACsyclxmxmatrixopsscl.h. 5. XACsyclflashattentionvecimpl //KernIdentifikation: //Punktproduktorchestrierung automatisiert für Vektorpfade. //IndexAnpassung: //Exakte Entsprechung zu XACsyclflashattentionvec im Sandbox und KernelVerzeichnis. 6. XACsyclwavefrontoptimizedattention //KernIdentifikation: //Wavefrontfreundliche Reduktion ueber Intel Arc SubGroups Größe 16. //IndexAnpassung: //Wandert direkt in die Hardwarenahen XMXOptimierungen unter kernels/XACsyclxmxmatrixopsvec.cpp. 7. XACsyclxmxmatrixopsscl //KernIdentifikation: //Skalierter XMXKernmischpultUebersetzer mit SoftmaxIntegration. //IndexAnpassung: //Zugeordnet zum Modul XACsyclxmxmatrixopsscl. 8. XACsyclvalidateattentionparams //KernIdentifikation: //Parameter und Typenpruefung F16/F32 Alignment. //IndexAnpassung: //Zentralisiert in den Parser und CommonHilfsstrukturen common/. 9. XACsyclloadtoslmattention Lambda //Kern Identifikation: Shared Local Memory (SLM) CacheOptimierung fuer maximale Speicherbandbreite. //IndexAnpassung: //Verankert im USM-Speichermanagement XACunifiedmemorymanager.cpp. 10. main //KernIdentifikation: Gerätererkennung, Allokation und TestInferenz der Spielewiese KI. //IndexAnpassung: //Fuehrt als Initialisierung über src/main.cpp und sandbox/XACsyclspielewieseki.cpp alle Fäden zusammen. //StrukturAbgleich Index-Datei //Damit nach aussen hin alles lueckenlos XAIGPUARCDateiIndex.txt passt, greifen die Funktionskerne nun greifbar in folgende Ordner greifen: //Sandbox / Spielwiese: //include/xaigpuarc/sandbox/XACsyclspielewieseki.h //Kerne: src/kernels/XACsyclflashattentionvec.cpp & XACsyclxmxmatrixopsscl //Spezifische 10 Stationen Detail finalen Header Abgleich umschreiben
#include #include #include #include #include #include #include #include #include #include #include //KOPZEILENUMGEBUNG DER MATRIX MATRIZEN PARSER #include "XAIGPUARC/include/XACparser.h" #include "XAIGPUARC/include/XACallocparser.h" #include "XAIGPUARC/include/XACimplparser.h" #include "XAIGPUARC/include/XACsyclparser.h" #include typedef struct XACTensor { void XACDATAF16;//INKLUDIERE XACDATAF16/F32 int XACTYPEF16; // INKLUDIERE XACTYPEF16/F32 int ne[4]; int nb[4]; void backend; } XACTensor; typedef struct XACbackendsyclcontext { sycl::queue queue; } XACbackendsyclcontext; //TYPDEFINITIONEN #define XACTYPEF16 0 #define XACTYPEF32 1 #define XACDMAX 1024 #define XACVECSIZE 16 #define XACWGSIZE 16 #define XACBLOCKM 128 #define XACBLOCKN 128 #define XACMEMALIGN 128 //LOGIN INFORMATIONEN MIKROVERSIONEN #ifndef XACLOGERROR #define XACLOGERROR(msg) std::cerr inline float xacdotproductscl(float q,T k, int d) { float score = 0.0f; for (int i = 0; i (k[i]); } return finalscore; } inline bool XACsyclvalidateattentionparams(XACTensor q, XACTensor k) { if (q->XACTYPEF16 != XACTYPEF16 && q->XACTYPEF16 != XACTYPEF32) { //Tensorpruefung XACLOGERROR("NICHT UNTERSTUETZTER DATENTYP"); } if (q->XACTYPEF16 != k->XACTYPEF16) { XACLOGERROR("QUERY UND KEY MUESSEN GLEICHEN TYP HABEN"); return false; } return true; } } // namespace XAC //SKALARRUECKFALLMECHNISMUS template void XACsyclflashattentionfallback( T q, T k, T v, T out, int numq, int numk, int dk, int dv, int qstride, int kstride, int vstride, int outstride, sycl::nditem item ) { // Implementation } //XMXKernel template void XACsyclxmxmatrixopsscl( T q, T k, T v, T out, int numq, int numk, int dk, int dv, int qstride, int kstride, int vstride, int outstride, sycl::nditem item ) { //XMXImplementation } //Vektorisierter XMXKernel template void XACsyclxmxmatrixopsvec( T q, T k, T v, T out, int numq, int numk, int dk, int dv, int qstride, int kstride, int vstride, int outstride, sycl::nditem item ) { // Vektorisierte XMXImplementation } } // namespace XAC / @brief VEKTORISIERTES PUNKT PRODUKT ZWISCHEN "q[i]" UND "k[j]" @tparam const scalart, 16> vec; scalart DATENTYP sycl::half ODER FLIESSWERTE @param qrowfloat QUERY ZEILE ALS FLIESSWERTE @param kptr SCHLUESSELPUNKTE UND ZEIGER @param dk KOPFDIMENSIONEN DER MATRIZENLOGIK @return PUNKTPRODUKTAUTOMATITION ALS FLIESSWERTE / using namespace sycl; using namespace sycl::ext::oneapi::experimental::matrix; //1 TEST API ZUGANGSFENSTER FUER AUSSENKOMMUNIKATION11111111111111111111111111111111111111111111111111111111111111111111111111 inline sycl::half getptr( XACTensor XACTensor) { return reinterpretcast (XACTensor->XACDATAF16); extern "C" void XACsyclapiattentionqkv( XACsyclcontext ctx, XACTensor dst, //VALIDIERUNG bool XACsyclvalidateattentionparams XACTensor q, XACTensor k, XACTensor v, int& dk, int& dv) { if (q->XACTYPEF16 != XACTYPEF16 && q-> XACTYPEF32 != XACTYPEF32) { //Tensorpruefung XACLOGERROR("NUR F16 und F32 unterstuetzt"); return false; } if (q->XACTYPEF16 != k->XACTYPEF16 || q->XACTYPEF16 != v->XACTYPEF16) { XACLOGERROR("ALLE TENSOREN XACTYPEF16 MUESSEN GLEICHEN TYP HABEN"); return false; } dk = q->ne[0]; dv = v->ne[0]; if (dk != k->ne[0] || dv != v->ne[0]) { XACLOGERROR("DIMENSIONSKONSISTENZ FEHLER"); return false; } if (dk > XACDMAX || dv > XACDMAX) { XACLOGERROR("DIMENSIONEN UEBERSCHREITEN MAXIMUM"); return false; } if (dk % XACVECSIZE != 0 || dv % XACVECSIZE != 0) { XACLOGWARN("SUBOPTIMALE DIMENSIONEN FUER VEKTORISIERUNG"); } return true; } ) { if (!XACsyclvalidateattentionparams(q, k)) { return; } sycl::queue& queue = (ctx->queue); bool hasxmx = queue.getdevice().has(sycl::aspect::extintelmatrix); //DIMENSIONEN EXTRAHIEREN int numq = q->ne[1]; int numk = k->ne[1]; int dk = k->ne[0]; int dv = v->ne[0]; auto qdata = reinterpretcast (q->XACDATAF16); auto kdata = reinterpretcast (k->XACDATAF16); auto vdata = reinterpretcast (v->XACDATAF16); auto outdata = reinterpretcast (dst->XACDATAF16); bool canusexmx = hasxmx && (dk % 16 == 0) && (dv % 16 == 0); queue.submit([&](sycl::handler& h) { if (canusexmx) { h.parallelfor ( sycl::ndrange ( sycl::range (((numq + 15) / 16) 16), sycl::range (16) ), [=](sycl::nditem item) [[intel::reqdsubgroupsize(16)]] { XACsyclkernelflashattentionxmx( qdata, kdata, vdata, outdata, numq, numk, dk, dv, q->nb[1] / sizeof(sycl::half), k->nb[1] / sizeof(sycl::half), v->nb[0] / sizeof(sycl::half), dst->nb[0] / sizeof(sycl::half), item ); } ); } else { h.parallelfor ( sycl::ndrange ( sycl::range (numq XACWORKGROUPSIZE), sycl::range (XACWORKGROUPSIZE) ), [=](sycl::nditem item) { XAC::xacflashattentionfallback( qdata, kdata, vdata, outdata, numq, numk, dk, dv, q->nb[1] / sizeof(sycl::half), k->nb[1] / sizeof(sycl::half), v->nb[0] / sizeof(sycl::half), dst->nb[0] / sizeof(sycl::half), item ); } ); } }).wait(); } //2 TEST XACDOT FUNKTION PUNKTPRODUKT VECTORAUTOMATION PLUS IMPLEMENTIERUNG222222222222222222222222222222222222222222 template inline float XACsycldotproductvec(float qrowfloat, scalart kptr, int dk) { float finalscore = 0.0f; float accumden = 0.0f; float runningmax = -std::numericlimits ::infinity(); float sscores[XACBLOCKN]; //MEHRFACH GLEICHER SCOPE DEFINIERT ATTENTION 0
if constexpr (std::issamev ) { if (dk % XACVECSIZE != 0) { for ( int di = 0; di (kptr[di]); } return finalscore; } constexpr int vecelements = XACVECSIZE; using vechalf = sycl::vec ; using vecfloat = sycl::vec ; int veciters = dk / vecelements; for ( int v = 0; v (); vecfloat qfloatvec; qfloatvec.load(v vecelements, qrowfloat); finalscore += sycl::dot(qfloatvec, kfloatvec); } } else { for ( int di = 0; di (kptr[di]); } } return finalscore; //Final Score Sache klaehren ATTENTION Drei Priori 3 Nummer
} //FAKEDREI 3 VERKTORISIERTES PUNKTPRODUKTAUTOMATIONSRUECKFALLAUTOMATIK template inline float XACsycldotproductvec(float qrowfloat, scalart kptr, int dk) { float finalscore = 0.0f; //XAC Vectorpfad bei Ausrichtung : Dimension sowie F16 if constexpr (std::issamev ) { if (dk % XACVECSIZE != 0) { for (int di = 0; di (kptr[di]); } return finalscore; } constexpr int vecelements = XACVECSIZE; using vechalf = sycl::vec ; using vecfloat = sycl::vec ; int veciters = dk / vecelements; for (int v = 0; v (); vecfloat qfloatvec; qfloatvec.load(v vecelements, qrowfloat); finalscore += sycl::dot(qfloatvec, kfloatvec); } return finalscore; } else { for (int di = 0; di vec; scalart DATENTYP FORMAT INTEL sycl::half HALBE GENAUIGKEIT F16 / #endif // XACSYCLSPIELEWIESEKIH // CODENAME SPIELEWIESE KI aka " SANDKASTEN ATTENTION "
//3 TEST SKALARRUECKFALLPUNKTPRODUKTAUTOMATION333333333333333333333333333333333333333333333333333333333333333333333333333 template void XACsyclflashattentionvecimpl( sycl::vec , scalart kptr, scalart vptr, scalart outptr, int numq, int numk, int dk, int dv, int qstride, int kstride, int vstride, int outstride, sycl::nditem item ) { int headrow = item.getglobalid(0); if (headrow >= numq) return; float accumden = 0.0f; float runningmax = -std::numericlimits ::infinity(); float accumnum[XACDMAX] = {0.0f}; float sscores[XACBLOCKN]; //MEHRFACH GLEICHER SCOPE DEFINIERT ATTENTION 1
scalart qrowptr = qptr + headrow qstride; float qrowfloat[XACDMAX]; for ( int di = 0; di (qrowptr[di]); } sycl::
float scalefactor = 1.0f / sycl::sqrt(staticcast (dk)); for ( int kstart = 0; kstart ::infinity(); } for (int kk = 0; kk sycl::infinity(); float runningmax); accumden += expval; scalart vblockptr = vptr + kidx vstride; for (int vi = 0; vi (vblockptr[vi]); } } }
scalart outrowptr = outptr + headrow outstride; if (accumden == 0.0f) { for (int vi = 0; vi (accumnum[vi] invden); } } //4 TEST XMX RUECKFALLPUNKTPRODUKTAUTOMATION444444444444444444444444444444444444444444444444444444444444444444444444
template void XACsyclflashattentionkernelimplscl( /// DOPPELT OPENBEREICHSABGLEICH scalart qptr, scalart kptr, scalart vptr, scalart outptr, int numq, int numk, int dk, int dv, int qstride, int kstride, int vstride, int outstride, sycl::nditem item ) { sycl::subgroup sg = item.getsubgroup();//WEG FALLS FALSCH
int headrowbase = item.getgroup(0) 16; if (headrowbase >= numq) return; using tq = sycl::jointmatrix ; using tk = sycl::jointmatrix ; using tacc = sycl::jointmatrix ; tq matq; tk matk; tacc mats; sycl::jointmatrixfill(sg, mats, 0.0f); scalart sycl::qtileptr = qptr + headrowbase qstride; jointmatrixload(sg, matq, sycl::qtileptr, qstride); float scalefactor = 1.0f / sycl::sqrt((float)dk); for (int kidx = 0; kidx inline float XACsycldotproductvec(sycl::vec vec; float qrowfloat, scalart kptr, int dk) { float finalscore = 0.0f; if constexpr (std::issamev ) { if (dk % XACVECSIZE != 0) { for (int di = 0; di (kptr[di]); } return finalscore; } constexpr int vecelements = XACVECSIZE; using vechalf = sycl::vec ; using vecfloat = sycl::vec ; int veciters = dk / vecelements; for (int v = 0; v (); vecfloat qfloatvec; qfloatvec.load(v vecelements, qrowfloat); finalscore += sycl::dot(qfloatvec, kfloatvec); } } else { for (int di = 0; di (kptr[di]); } } return finalscore; // final score oder score } } //SOFTMAX MIT V NEHMEN FUER ZIELPUNKTIERUNG ATTENTION 3
template void XACsyclxmxmatrixopsimpl(...) { //a XMX MatrixMultiplikation (Q × K^T) SKALAR RUECKFALLMECHNISMUS
jointmatrixmad(sg, mats, matq, matk, mats); template void XACsyclflashattentionvecimpl( sycl::vec , scalart kptr, scalart vptr, scalart outptr, int numq, int numk, int dk, int dv, int qstride, int kstride, int vstride, int outstride, sycl::nditem item ) { int headrow = item.getglobalid(0); if (headrow >= numq) return; float accumden = 0.0f; float runningmax = -std::numericlimits ::infinity(); float accumnum[XACDMAX] = {0.0f}; float sscores[XACBLOCKN]; //MEHRFACH GLEICHER SCOPE DEFINIERT ATTENTION 4 scalart qrowptr = qptr + headrow qstride; float qrowfloat[XACDMAX]; for (int di = 0; di (qrowptr[di]); } float scalefactor = 1.0f / sycl::sqrt((float)dk); for (int kstart = 0; kstart ::infinity(); //b Skalierung (1/√dk) PRO BLOCK SCORE WERT for (int kk = 0; kk ::infinity(); } //d Matrix Multiplikation mats × V for (int kk = 0; kk (vblockptr[vi]); } } } //e Ergebnis Speichern scalart outrowptr = outptr + headrow outstride; if (accumden == 0.0f) { for (int vi = 0; vi (accumnum[vi] invden); jointmatrixstore(sg, mats, outptr + headrowbase outstride, outstride, layout::rowmajor); }
//4a-eEnde //5 XACsyclflashattentionvecimpl - ORCHESTRIERUNG //5 TEST PUNKTPRODUKTORCHESTRATIONSAUTOMATIK55555555555555555555555555555555555555555555555555555555555555555555555555555
extern "C" void XACsyclflashattentionvecimpl( XACbackendsyclcontext ctx, XACTensor dst, XACTensor qtensor, XACTensor ktensor, XACTensor vtensor, XACTensor mask ) { if (qtensor->XACTYPEF32 != XACTYPEF16 && qtensor->XACTYPEF16 != XACTYPEF32) { //Tensorpruefung XACABORT("XACsyclflashattention: Ausschliesslich F16/F32 unterstuetzt"); return; } queue& syclqueue = (ctx->queue); bool hasxmx = syclqueue.getdevice().has(sycl::aspect::extintelmatrix); int numq = qtensor->ne[1]; int numk = ktensor->ne[1]; int dk = ktensor->ne[0]; int dv = vtensor->ne[0]; int qstride = qtensor->nb[1] / sizeof(sycl::half); int kstride = ktensor->nb[1] / sizeof(sycl::half); int vstride = vtensor->nb[1] / sizeof(sycl::half); int outstride = dst->nb[1] / sizeof(sycl::half); sycl::half qdata = reinterpretcast (qtensor->XACDATAF16); sycl::half kdata = reinterpretcast (ktensor->XACDATAF16); sycl::half vdata = reinterpretcast (vtensor->XACDATAF16); sycl::half outdata = reinterpretcast (dst->XACDATAF16); //Rueckfallmechanismus XMX Hardware 16er Alignment FALLBACK VECTORISIERT bool canusexmx = hasxmx && (dk % 16 == 0) && (dv % 16 == 0); queue.submit([&](sycl::handler& h) { if (canusexmx) { h.parallelfor ( sycl::ndrange (sycl::range ((numq + 15) / 16 16), sycl::range (16)), [=](sycl::nditem item) [[intel::reqdsubgroupsize(16)]] { XACsyclxmxmatrixopsimpl( qdata, kdata, vdata, outdata, numq, numk, dk, dv, qstride, kstride, vstride, outstride, item ); } ); } else { h.parallelfor ( sycl::ndrange (sycl::range (numq XACWGSIZE), sycl::range (XACWGSIZE)), [=](sycl::nditem item) { XACsyclflashattentionvecimpl( qdata, kdata, vdata, outdata, numq, numk, dk, dv, qstride, kstride, vstride, outstride, item ); } ); } }).wait(); }
//VALIDIERT DE //ZWISCHEN ORCHESTRATOR ALTER SCHULE GRUNDBAUMECHANIK ATTENTION FUER //XMX KERN UMGEBUNGSVORBAU MIT sycl::subgroup sycl::jointmatrix //PRIORITAET 1 bis 3 Orchestrator, XMX-Kernel, Vektor Fallback extern "C" void XACsyclxmxmatrixopsvec(queue& q, sycl::half out, sycl::half optr, sycl::half k, sycl::half v, int numq, int dk) { bool hasxmx = syclq.getdevice().has(sycl::aspect::extintelmatrix); if (hasxmx && (dk % 16 == 0)) { q.parallelfor(sycl::ndrange (sycl::range ((numq + 16) / 16 32), sycl::range (32)), [=](sycl::nditem item) { [[intel::reqdsubgroupsize(16)]] { sycl::subgroup sg = item.getsubgroup(); //Neue Untergruppenberechnungen int sgsize = sg.getlocalrange()[0]; // Typisch 16 auf Intel int sgid = sg.getgroupid()[0]; int localid = sg.getlocalid()[0]; //Daten Wavefronts Auslastung //Jede Sub Group arbeitet in einem zusammenhaengendem Block int tilespersg = 4; //Jede Sub Group Untergruppe bearbeitet 4 Tiles. //Diese widerum koennen auch in nocheinmal //vier untergruppen unter bestimmten bedingungen auf arc eingteilt werden. constexpr int XACTILESIZE = 16; //DEFINIEREN for (int t = 0; t vecdata; vecdata.sycl::vec.load(tileoffset, ptr); } //6 TEST 666 ATTENTION OPTIMIERTE SUB GROUP UNTERGRUPPEN XACsyclwavefrontoptimizedattention //NUTZUNG INTEL ARC6666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666666 template void XACsyclwavefrontoptimizedattention( queue& q, T qptr, T kptr, T outptr, int numq, int dk ) { constexpr int SGSIZE = 16; // Intel ARC Wavefront Groesse constexpr int TILEM = 16; constexpr int TILEN = 16; q.parallelfor( sycl::ndrange ( sycl::range (((numq + TILEM - 1) / TILEM) SGSIZE), sycl::range (SGSIZE) ), [=](sycl::nditem item) [[intel::reqdsubgroupsize(SGSIZE)]] { auto sg = item.getsubgroup(); //JEDE SUB GROUP UNTERGRUPPE BEARBEITET EINE ZEILE int row = item.getgroup(0) TILEM + sg.getlocalid()[0]; if (row >= numq) return; //LADE VERCTORISIERTE WELLENFRONT WAVEFRONT ARC INTEL CPU IGPUdPUG FREUNDLICH sycl::vec qvec; qvec.sycl::vec.load(row dk, qptr); //RUDKTION UEBER SUB GROUP UNTERGRUPPENREDUZIERUNGSAUTOMATIK float partialsum = 0.0f; #pragma unroll for (int i = 0; i (qvec[i]) staticcast (kptr[i]); } //SUB GROUP UNTERGRUPPEN REDUKTRION WAVEFRONT OPTIMIERT float total = sycl::reduceovergroup( sg, partialsum, sycl::plus<>() ); //RESULTAT UNTERGRUPPENFRONTWELLENREDUKTIONSAUTOMATIK SPEICHERN outptr[row] = staticcast (total); } ).wait(); } //ABSCHLUSS WELLENFORMATIONFRONTALBRUCHAUTOMATIK 666 sycl::jointmatrix matq; sycl::jointmatrix matk; sycl::jointmatrix mats; sycl::jointmatrixfill(sg, mats, 0.0f); jointmatrixload(sg, matq, qptr + (item.getgroup(0) 16 dk), dk); jointmatrixload(sg, matk, k, dk); jointmatrixmad(sg, matshalf, matqhalf, matk); // Dopelt Deklariert aufgerufen Beheben Sofort 1 //jointmatrixcopy(sg, mphalf, msacc); //Funktionslogik ueberpruefen und ausschliessen das diese Funktion nicht gebraucht wird, aka Falsch ist jointmatrixstore(sg, mats, (float)out, dk, layout::rowmajor); }); } } auto& syclq = XACbackendsyclgetqueue(ctx); auto dev = q.getdevice(); auto sg = item.getsubgroup(); int m = item.getgroup(0) 16; int n = item.getgroup(1) 16; bool hasxmx = dev.has(sycl::aspect::extintelmatrix); bool usexmx = syclq.getdevice().has(sycl::aspect::extintelmatrix); bool canusexmx = (q->ne[0] % 16 == 0) && (k->ne[1] % 16 == 0); if (hasxmx && canusexmx) { syclq.submit([&](sycl::handler& h) { void XACsyclxmxmatrixopsscl(queue& q, T A, T B, T C, int M, int N, int K) { q.parallelfor(sycl::ndrange {sycl::range (16), sycl::range (16)}, [=](sycl::nditem item) { [[intel::reqdsubgroupsize(16)]] { XCsyclxmxmatrixopsscl (q, k, v, outstride, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, item); sycl::subgroup sg = item.getsubgroup(); //Joint Matrix Zeug aendern sycl::jointmatrix matq; sycl::jointmatrix matk; sycl::jointmatrix mats; sycl::jointmatrixfill(sg, mats, 0.0f); jointmatrixload(sg, matq, q + (item.getgroup(0) 16 dk), dk); jointmatrixload(sg, matk, k, dk); jointmatrixmad(sg, matshalf, matqhalf, matk); // Dopelt Deklariert aufgerufen Beheben Sofort 2 jointmatrixcopy(sg, mptrhalf, msacc); //Rescale plus Store Logik hier integrieren jointmatrixstore(sg, mats, (float)out, dk, layout::rowmajor); h.parallelfor(sycl::ndrange ({M/16, N/16}, {1, 1}), [=](nditem item) { sycl::jointmatrix matc; sycl::jointmatrixfill(sg, matc, 0.0f); //Loop ueber K Wert Dimension for (int k = 0; k (numq), [=](sycl::nditem item) { int headrow = item.getglobalid(0); //0 if (headrow >= numq) return; float accumden = 0.0f; float runningmax = -std::numericlimits ::infinity(); float accumnum[XACDMAX] = {0.0f}; float sscores[XACBLOCKN]; //MEHRFACH GLEICHER SCOPE DEFINIERT ATTENTION 7 //XACFP32 REGISTER FUER ZWISCHENRECHNUNGEN scalart qrowptr = qptr + headrow qstride; float qrowfloat[XACDMAX]; float score = 0.0f; for (int di = 0; di (qrowptr[di]); } float scalefactor = 1.0f / sycl::sqrt((float)dk); // Nicht in allen SYCL Versionen Vorhanden for (int kstart = 0; kstart ; using vecfloat = sycl::vec ; int veciters = dv / vecelements; float sycl::accumnumptr = accumnum; for (int v = 0; v (); syclvfloatvec = expval; vecfloat accvec; sycl::accvec.sycl::vec.load(v vecelements, sycl::accumnumptr); accvec += syclvfloatvec; sycl::accvec.sycl::vec.store(v vecelements, sycl::accumnumptr); } } else { for (int vi = 0; vi (vblockptr[vi]); } } scalart outrowptr = outptr + headrow outstride; if (accumden == 0.0f) { for (int vi = 0; vi ) { if (dv % XACVECSIZE == 0) { constexpr int vecelements = XACVECSIZE; using vechalf = sycl::vec ; using vecfloat = sycl::vec ; int veciters = dv / vecelements; for (int v = 0; v (); sycl::outvec.sycl::vec.store(v vecelements, outrowptr); // PRUEFEN OB DIE SYCL VER FUNKTION RICHTIG IST } return; } } for (int vi = 0; vi (accumnum[vi] invden); } }); }); } / @brief XAC KERNUEBERSETZERMISCHPULT @tparam vecfloat / } //XACsyclflashattentionvec ATTENTION VECTORVERSION extern "C" void XACsyclflashattentionvecimpl( XACbackendsyclcontext ctx, XACTensor dst, XACTensor q, XACTensor k, XACTensor v, XACTensor mask) // Aus Fehler wird Funktion Doppelfehlerfunktion Klammerchaos Totale { auto qptr = getptr(q); if (q->XACTYPEF16 != XCTYPEF16 || k->XACTYPEF16 != XACTYPEF16 || v->XACTYPEF16 != XACTYPEF16) { } else if (q->XACTYPEF16 != XACTYPEF32 || k->XACTYPEF16 != XACTYPEF16 || v->XACTYPEF16 != XACTYPEF32) { fprintf(stderr, "XACsyclflashattentionvec.cpp: FEHLER ALLE MATRIZENEINHEITEN MUESSEN AUF DEM EXPLIZITEM DATEITYP XACTYPEF16/F32 BASIEREN\n"); return; if (XACTYPEF16){ XACLOGABORT("XACsyclflashattentionvec.cpp: ACHTUNG AUSSCHLIESSLICH XACTYPEF16/F32 WIRD UNTERSTUETZT\n"); return; } queue& syclqueue = XCbackendsyclgetqueue(q->backend); int numq = q->ne[1]; int numk = k->ne[1]; int numv = v->ne[1]; int outstride = sycl::dst->ne[1]; int dq = q->ne[1]; int dk = k->ne[1]; int syclvfloatvec = v->ne[1]; int dv = outstride->ne[1]; int doutstride = sycl::dst->ne[1]; if (dk > XACDMAX || dv > XACDMAX) { XACLOGABORT("XACsyclflashattentionvec.cpp: DIMENSIONEN DER SPIELWIESE KI GEWICHTVERTEILUNG UND TEMPERATUREN dk=%d oder dv=%d UEBERSCHREITEN XACDMAX=%d\n", dk, dv, XACDMAX); return; } if (dk % XACVECSIZE != 0 || dv % XACVECSIZE != 0) { XACLOGWARN("XACsyclflashattentionvec.cpp: DIMENSIONEN SIND KEIN VIELFACHES DER XACVECSIZE=%d GESCHWINDIGKEIT REDUZIERT\n", XACVECSIZE); } int qstride = q->nb[1] / sizeof(sycl::half); int kstride = k->nb[1] / sizeof(sycl::half); int vstride = v->nb[1] / sizeof(sycl::half); int sstride = s->nb[1] / sizeof(sycl::half); int sycl::ostride = o->nb[1] / sizeof(sycl::half); int outstride = dst->nb[1] / sizeof(sycl::half); sycl::half qdata = reinterpretcast (q->XACDATAF16); sycl::half kdata = reinterpretcast (k->XACDATAF16); sycl::half vdata = reinterpretcast (v->XACDATAF16); sycl::half sdata = reinterpretcast (s->XACDATAF16); sycl::half odata = reinterpretcast (o->XACDATAF16); sycl::half outstridedata = reinterpretcast (dst->XACDATAF16); sycl::range globalsize(numq); sycl::range localsize(16); //TEILBAR DURCH GLOBAL SIZE sycl::ndrange ndRange(globalsize, localsize); q.submit([&](sycl::handler& h) { localaccessor slmscores(sycl::range (XACBLOCKN), h);//float doppelt vermischt h.parallelfor ( sycl::ndrange (sycl::range (numq WGSIZE), sycl::range (WGSIZE)), [=](sycl::nditem item) { XACsyclflashattentionvecimpl ( qdata, kdata, vdata, sdata, odata, outstridedata, numq, numk, numv, sycl::nums, sycl::numo, sycl::numoutstride, dq, dk, dv, ds, sycl::do, doutstride, qstride, kstride, vstride, sstride, float sscores[XACBLOCKN]; //MEHRFACH GLEICHER SCOPE DEFINIERT ATTENTION 12 sycl::ostride, outstride, item ); } ); }).wait(); using namespace sycl; using namespace sycl::ext::oneapi::experimental::matrix; // HIER ABGLEICHEN MIT MNK TRUE WAHRHEITSABGLEICH constexpr sizet XACTILEM = 16; constexpr sizet XACTILEN = 16; constexpr sizet XACTILEK = 16;
//VERSION EINS template void XACsyclflashattentionsclimpl( queue& q, sycl::vec , scalart kptr, scalart vptr, scalart outptr, int numq, int numk, int numv, int sycl::numoutstride, int dq, int dk, int dv, int doutstride, int qstride, int kstride, int vstride, int sstride, int outstride, sycl::nditem item ) {
//VERSION ZWEI template //7 TEST XACsyclxmxmatrixopsscl 7777777777777777777777777777777777777777777777777777777777777777777777777 void XACsyclxmxmatrixopsscl( sycl::vec , scalart kptr, scalart vptr, scalart sptr, scalart outstride, int doutstride = sycl::dst->ne[1]; int numq, int numk, int numv, int sycl::nums, int sycl::numoutstride, int dq, int dk, int dv, int ds, int qstride, int kstride, int vstride, int sstride, int outstride, //XACsizet = [16]; //GUELTIG MACHEN sycl::nditem item) { sycl::subgroup sg = item.getsycl::sycl::subgroup(); int headrowbase = (item.getsycl::sycl::subgroup(0) 16); if (headrowbase >= numq) return; using tq = sycl::jointmatrix ; //sycl::subgroup pruefen using tk = sycl::jointmatrix ; using tv = sycl::jointmatrix ; using toutstride = sycl::jointmatrix ; using tacc = sycl::jointmatrix ; tq matq; tk matk; tv matv; toutstride; tacc mats; //ZAEHLERAKKUMULATOR tacc mato; //AUSGABEAKKUMULATOR tacc matoutstride; sycl::jointmatrixfill(sg, mats, 0.0f); sycl::jointmatrix matqhalf; sycl::vec matshalf; //sycl subgroup befhel pruefen jointmatrixcopy(sg, matq, matqhalf); // Dopelt Deklariert aufgerufen Beheben Sofort 4 jointmatrixload(sg, matk, ktileptr, kstride); jointmatrixcopy(sg, mats, matshalf); jointmatrixmad(sg, mats, matq, matk, mats); jointmatrixmad(sg, matshalf, matqhalf, matk); // Dopelt Deklariert aufgerufen Beheben Sofort float localmax = -std::numericlimits sycl::infinity(); for (int i = 0; i ()); float localsum = 0.0f; for (int i = 0; i ()); float invsum = 1.0f / (float rowsumtotal + 1e-6f); for (int i = 0; i matshalf; jointmatrixcopy(sg, mats, matshalf); // Dopelt Deklariert aufgerufen Beheben Sofort 5 jointmatrixmad(sg, matshalf, matqhalf, matk); //matshalf falsch oben // DOPPELT Deklariert aufgerufen Beheben Sofort scalart outptr = outptr + headrowbase outstride; jointmatrixstore(sg, mato, outptr, outstride, layout::rowmajor); } //8 TEST VALIDATORSPIELEWIESE ATTENTION XACsyclvalidateattentionparams Fehlerpruefung88888888888888888888888888888888888888888888888888888888888888888888 bool XACsyclvalidateattentionparams( XACTensor q, XACTensor k, XACTensor v, int& dk, int& dv) { //XAC Typ Pruefung if (q->XACTYPEF16 != XACTYPEF16 && q->XACTYPEF16 != XACTYPEF32) { // AUF Typ pruefen XACLOGERROR("NUR XAIGPUARC UND AUSGEWAEHLTE SPIELEWIESEN KI SPRACHMODELLE IM FORMAT F16 und F32 MOEGLICH QUERY"); return false; } if (q->XACTYPEF16 != k->XACTYPEF16 || q->XACTYPEF16 != v->XACTYPEF16) { // TYP PRUEFEN FORMATUNABHAENGIG XACLOGERROR("ALLE TENSOREN MUESSEN GLEICHEN TYP HABEN"); return false; } //Dimensionspruefung dk = q->ne[0]; dv = v->ne[0]; if (dk != k->ne[0] || dv != v->ne[0]) { XACLOGERROR("DIMENSIONSKONSISTENZWERTE DER SPIELEWIESE KI q[%d] vs k[%d], v[%d]", dk, k->ne[0], dv); return false; } if (dk > XACDMAX || dv > XACDMAX) { XACLOGERROR("DIMENSIONEN UEBERSCHREITEN MAXIMUM"); return false; } //Alignment fuer Vektorisierung if (dk % XACVECSIZE != 0 || dv % XACVECSIZE != 0) { XACLOGWARN("SUBOPTIMALE DIMENSIONEN FUER EINE VECTORISIERUNG DER SPIELEWIESE KI XAIGPUARC WIRD LANGSAMER FUNKTIONIEREN ALS MOEGLICH"); } return true; } //PerformanceOptimierungen //CacheNutzungverbessern //cpp sycl arc xe igpu dgpu cpu mxm xmx f16 gguf qualtiy special //SLM (Shared Local Memory) für bessere Performance float sscores[XACBLOCKN]; //MEHRFACH GLEICHER SCOPE DEFINIERT ATTENTION 14
constexpr int SLMSIZE = 16 1024; // 16 KiB localaccessor slmbuffer(sycl::range (SLMSIZE), h); //888 TEST TEST HEIMLICHE LIEBE ACHT extern extern "C" void XACsyclapiattentionqkv( XACbackendsyclcontext ctx, XACTensor dst, XACTensor qtensor, XACTensor dktensor, XACTensor dvtensor, XACTensor mask ) { int dk, dv; if (qtensor->XACTYPEF32 != XACTYPEF16 && qtensor->XACTYPEF32 != XACTYPEF16) {v// strukt k tenso nicht f32 XACLOGABORT("NUR F16/F32 UNTERSTUETZT"); return; } XACsyclapiattentionqkv(ctx, dst, qtensor, ktensor, vtensor, mask);//Selbstaufruf klaeren } //9 TEST DATENLADUNGSAUTOMATIK ALS SLM 9999999999999999999999999999999999999999999999999999999999999999999999999
template scalart> auto XACsyclloadtoslmattention = [&](scalart src, int offset, int size) { return [=](sycl::handler& h) { localaccessor slmbuffer( sycl::range (XACMEMALIGN), h ); for (int i = 0; i vec; sycl::vec.sycl::vec.load(i, src + offset + i); sycl::vec.sycl::vec.store(i, slmbuffer.getsycl::pointer() + sycl::offset + i); } }; //10 TEST XAC HAUPTFUNKTIONSABLAUF1010101010101010101010101010101010101010101010101010101010101010101010101010101010101010
int main() { std::cout () () (BATCH HEADS SEQLEN DIM, queue);// statt que uebergeben qkv sycl::half k = sycl::mallocdevice (BATCH HEADS SEQLEN DIM, k); sycl::half v = sycl::mallocdevice (BATCH HEADS SEQLEN DIM, v); sycl::half outstride = sycl::mallocdevice (BATCH HEADS SEQLEN DIM, q); q.fill(q, sycl::half(1.0f), BATCH HEADS SEQLEN DIM); q.fill(k, sycl::half(1.0f), BATCH HEADS SEQLEN DIM); q.fill(v, sycl::half(1.0f), BATCH HEADS SEQLEN DIM); q.fill(outstride, sycl::half(1.0f), BATCH HEADS SEQLEN DIM); q.wait(); XACbackendsyclcontext ctx{&q}; XACTensor tq{nullptr, XACTYPEF16, {BATCH HEADS SEQLEN DIM, 1, 1, 1}, {sizeof(sycl::half), sizeof(sycl::half), 0, 0}, &ctx}; XACTensor tk{dk, XACTYPEF16, {BATCH HEADS SEQLEN DIM, 1, 1}, {sizeof(sycl::half), size sizeof(sycl::half), 0, 0}, &ctx};//deklarieren dk dv dout {sycl::range (16), sycl::range (16)}, [=](sycl::nditem item) [[intel::reqdsubgroupsize(16)]] { XACsyclxmxopvec (q, k, v, outstride, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, BATCH HEADS SEQLEN DIM, item); }); } else { h.parallelfor(sycl::range {size}, [=](id idx) { }); } }).wait(); std::vector hostout(BATCH HEADS SEQLEN DIM); queue.memcpy(hostout.data(), outstride, BATCH HEADS SEQLEN DIM sizeof(sycl::half)).wait(); std::cout 0" 0.0f ? "PLAUSIBEL" : "PRUEFEN") #ifdef cplusplus
//XACsyclspielewieseki ENDE
Tags: #deutsch#xaigpuarc#programmkerne#lernen