Im Rahmen der Kolumne "Eine Frage an einen Intel-Experten stellen" haben wir den führenden Intel-Spezialisten Konstantin Vladimirov gebeten, Fragen zu heterogener Programmierung, dem
oneAPI-Toolkit und ähnlichen interessanten Dingen zu beantworten. Das Ergebnis hat alle unsere Erwartungen übertroffen. Konstantin hatte keine freie Zeit und gab detaillierte und fundierte Antworten, ohne Polemik zu befürchten. Tatsächlich haben wir einen kleinen Vortrag über die Architekturübergreifende Programmierung in all ihren Formen gehalten: Nuancen, Optimierungen, Standards und so weiter auslagern.
Wir übergeben das Mikrofon dem Fachmann. Nun, die Kommentare werden dem Publikum gegeben.
Frage
Soarex16Wie mühsam wird der Übergang von OpenCL zu oneAPI sein und welche Vorteile lassen sich daraus ziehen?
Die antwort. Der Wechsel zu DPC ++ kann schwierig sein, aber meiner Meinung nach lohnt es sich. Es gibt zwei Hauptstufen.
Erstens ist dies ein Übergang von Ihrer heterogenen Programmiersprache (OpenCL, Vulkan Compute), die höchstwahrscheinlich auf der API basiert. Hier haben Sie einen Vorsprung, da Sie den Themenbereich bereits kennen und die Schwierigkeit darin besteht, das Denken von der direkten Steuerung über die API auf etwas implizitere Sprachkonstrukte umzustellen.
Zweitens ist dies ein Übergang von Ihrer Gastsprache. Wenn Sie Ihr ganzes Leben von reinem C abgeladen haben, entspricht der Eingabeschwellenwert dem Schwellenwert für den Wechsel von C zu C ++, der ziemlich hoch ist.
Warum versuchen?
Erstens leistet DPC ++ einen großartigen Job für einen Programmierer. Sie werden sehr schnell vergessen, wie ein Albtraum, all diese expliziten Aufrufe von clXXXYYY, und was das sechste Argument bedeutet, und ob Sie den Rückkehrcode vergessen haben. Viele objektorientierte Wrapper verstecken die Routine nicht schlechter, aber normalerweise auf Kosten der Umstellung von der Standard-OpenCL-API auf die nicht so Standard-Wrapper-API (ich habe diese Bikes auch gesehen). Im Falle von DPC ++ schreiben Sie einfach die Standard-SYCL mit Intel-Erweiterungen (die möglicherweise bald auch zur Standard-SYCL werden).
Zweitens sorgt DPC ++ für eine gemeinsame Kompilierung, dh Sie können sicher sein, welche Typen es gibt, und Sie haben keine Probleme an den Grenzen der API mit Bemaßungen, Abständen und Ausrichtungen. Sie schreiben den Kernel- und Host-Code in eine Datei, und dies ist derselbe Code. Mit USM können Sie auch viel einfacher mit komplexen Datenstrukturen arbeiten.
Drittens ist DPC ++ echtes C ++, das heißt, es ermöglicht eine verallgemeinerte Programmierung. Zum Beispiel der einfachste Kernel zum Hinzufügen von zwei Vektoren:
auto kern = [A, B, C](cl::sycl::id<1> wiID) { C[wiID] = A[wiID] + B[wiID];
Dasselbe auf OpenCL:
_kernel void vector_add(__global int *A, __global int *B, __global int *C) { int i = get_global_id(0); C[i] = A[i] + B[i]; }
Sie sehen, ich war gezwungen, auf einen OpenCL-Typ int zu zeigen. Wenn ich einen Float benötige, muss ich entweder einen anderen Kernel schreiben oder einen Präprozessor oder eine externe Codegenerierung verwenden. Es kann ein wenig beängstigend sein, fast alle Funktionen von C ++ zur Verfügung zu haben, wenn Sie noch keine Erfahrung mit C ++ haben. Dies ist jedoch eine häufige Angelegenheit, wenn es um einen großen technologischen Wandel geht.
Und alle Vorteile sind nicht darauf beschränkt. Ich werde in den folgenden Antworten etwas anderes erwähnen.
Daher hätte ich den Compiler an Ihrer Stelle heruntergeladen und ausprobiert, da dies mit dem
OneAPI- Paket nicht schwierig ist.
Frage
JusterWerden OpenVINO und oneAPI in irgendeiner Weise zusammenhängen?
Die antwort. Die OpenVINO-Distribution ist jetzt Teil der OneAPI-Distribution. Das Erlernen und Verwenden neuronaler Netze ist eine rechenintensive Aufgabe, die stark von heterogener Programmierung profitiert. Ich glaube, dass alle OneAPI-Komponenten es früher oder später ermöglichen werden, alle verfügbaren Rechenressourcen zu nutzen: sowohl Grafikbeschleuniger als auch Spezialbeschleuniger wie Nervana und FPGA. Und das alles, ohne das Sprachparadigma und Typensystem Ihres C ++ - Programms zu verlassen.
Fragen aus der Mail
Ich versuche zu verstehen, wie AI Hardware Accelerator in 3 Jahren aussehen wird. Bitte helfen Sie dabei. Es gibt ein interessantes Unternehmen Graphcore und seine IPU - dieses Gerät ist nicht weniger effizient als FPGA, aber es ist viel einfacher zu programmieren - Python mit Unterstützung für TensorFlow und andere Frameworks. Es stellt sich heraus, dass bei Erfüllung der Versprechen von Graphcore keine FPGAs auf dem Markt für maschinelles Lernen erforderlich sind. Python ist für Datenwissenschaftler viel praktischer als C ++.
Stimmen Sie zu, dass FPGA im Vergleich zu programmierbaren Python-Lösungen nicht für den Markt für maschinelles Lernen geeignet ist? Welche anderen weit verbreiteten FPGA-Anwendungen sehen Sie, wenn der ML-Markt verloren geht?
In welchen Anwendungen sehen Sie den unvermeidlichen Bedarf an heterogener Programmierung, wo Sie mit bequemeren Tools wie Python nicht auskommen können?
Die antwort. Ich warf einen kurzen Blick auf die Art der IPU. Noch ein Stück Eisen, auf das sich jeder entladen wird. Diese Typen konkurrieren mit der GPU und mit speziellen Beschleunigern und nicht mit dem FPGA.
Bei Aufgaben, für die ein spezielles Hardwareteil eingesperrt ist, schlägt es immer das FPGA, z. B. ist das Rendern von Videos auf einer Grafikkarte usw. besser. Aber in der Welt (einschließlich der ML-Welt) gibt es viele Aufgaben, für die nichts Besonderes erfunden oder veröffentlicht wurde, und hier wird FPGA immer unverzichtbar sein. Zum Beispiel, weil es um den Preis geht und um billig zu sein, muss eine spezialisierte Hardware massiv sein.
Angenommen, die angegebene IPU ist wirklich cool. Dies wird die heterogene Programmierung nicht aufheben, im Gegenteil, das Vorhandensein eines solch hervorragenden Beschleunigers wird sie beflügeln. Und es wird OneAPI und DPC ++ einen Riesenvorsprung verschaffen, denn früher oder später wird jemand sagen: "Ich möchte sowohl Ihre IPU als auch meine GPU von einem Programm aus verwenden." Eher früh, weil es um heterogene Programmierung geht. Ihre Bedeutung ist die Verlagerung einer geeigneten Aufgabe auf ein geeignetes Gerät. Eine Aufgabe kann von überall kommen. Und dieses Gerät kann alles sein, es kann sogar dasselbe Gerät sein, auf dem das Programm ausgeführt wird. Wenn Sie beispielsweise den in ISPC geschriebenen Kernel auslagern und die Vektorfunktionen von Xeon maximal nutzen, können Sie ihn selbst auslagern und dennoch einen erheblichen Gewinn erzielen. Das Hauptkriterium hierbei ist die Leistung. Nun, es wird nie zu viel Produktivität auf dieser Welt geben. Selbst mit den besten Beschleunigern der Welt.
Was Python und seine Bequemlichkeit angeht ... Ich muss sofort zugeben, dass ich dynamisch getippte Sprachen nicht mag: Sie sind langsam und statt eines normalen Kompilierungsfehlers müssen Sie zwei Stunden warten, bevor Sie aufgrund des falschen Typs in die Laufzeit wechseln. Ich verstehe aber nicht, wie schlimm es ist, dieselben Auslagerungen unter Python durchzuführen. In OneAPI ist übrigens bereits Intel Distribution für Python enthalten, was für verschiedene Testberichte äußerst praktisch ist.
Das heißt, in der Traumwelt der Python-Liebhaber schreiben Sie ein Programm darauf und entladen es auf alle Beschleuniger, die Sie mit OneAPI finden können, und nicht auf eine Reihe von herstellerspezifischen Bibliotheken. Eine andere Sache ist, dass Sie bei diesem Ansatz die End-to-End-Eingabe verpassen und in die äußerst prekäre Welt der API-basierten Programmierung zurückkehren. Vielleicht wird die Entwicklung von DPC ++ die Community ermutigen, aktiver geeignete Tools wie C ++ zu verwenden.
Frage aus der Mail
Performance versus OpenCL. Luxus muss besteuert werden - d. H. Gemeinkosten. Gibt es irgendwelche Maße?
Die antwort. Im Internet finden Sie viele Messungen mit unterschiedlichen Ergebnissen, abhängig vom Compiler, der Aufgabe und der Qualität der Implementierung. Als persönliche Recherche habe ich an einfachen Aufgaben (SGEMM, DGEMM) auf meinem Laptop (integrierte Skylake-Grafik) gemessen und festgestellt, dass es bis jetzt einen gewissen Rückgang gibt (innerhalb von Prozent). Aber es scheint mir, dass dies eine Folge der Tatsache ist, dass dies alles bisher Beta ist.
Theoretisch sollte das Ergebnis Beschleunigung sein, nicht Verlangsamung, dh im Prinzip sollte all dieser Luxus einen negativen Wert haben. Alles dreht sich um den Compiler. Wenn Ihr Programm aus einer einzigen Quelle besteht und als einzelnes Programm verarbeitet wird, bietet der Compiler fantastische, unglaubliche Optimierungsmöglichkeiten: Erstellen von allgemeinen Codes, Umkehren von Schleifen, Neuanordnen von Codeabschnitten und alles andere, was der Compiler im API-basierten Ansatz einfach nicht kann früher oder später wird sie definitiv mit einem einzigen Quellmodell lernen.
Darüber hinaus wird DPC ++ negative Kosten in Bezug auf die Entwicklungszeit verursachen. Ein einfaches Beispiel sind SYCL-Accessoren, mit denen der Compiler bereits Ereignisse anordnet und asynchrone Warteschlangen verwaltet.
deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto C = bufferC.template get_access<sycl_write>(cgh); .... deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto D = bufferD.template get_access<sycl_write>(cgh);
Hier sieht der Compiler, dass beide Pakete nur A und B lesen und unabhängige Puffer C und D schreiben. Infolgedessen sieht er die Möglichkeit, sie parallel zu senden, wenn es genügend globale Größen gibt.
Natürlich kann ein pedantisch geschriebenes OpenCL-Programm dies genauso gut, aber die Entwicklungszeit, die mit einem nicht trivialen Kernel verbracht wird, ist nicht vergleichbar.
Frage aus der Mail
Sind alle Möglichkeiten zur Optimierung von OpenCL-Anwendungen für DPC ++ relevant? Was ist neu, um sie hinzuzufügen?
Die antwort. Ich würde sagen, dass die meisten der subtilen manuellen Optimierungen, die von Kernel-Autoren durchgeführt werden, vom Compiler vorgenommen werden können und sollten. Auf die gleiche Weise halte ich es beispielsweise für schädlich, einen Inline-Assembler manuell in C ++ - Programmen zu installieren, da er, selbst wenn er taktische Vorteile bietet, Optimierungen beeinträchtigt und sich negativ auf die Entwicklung und den Transfer eines Produkts auswirkt. Nun, OpenCL ist jetzt auch Assembler.
Was die detailliertere Antwort betrifft, habe ich Angst vor dem Abgrund hier. Beispielsweise gibt es ein bekanntes Intel-Dokument "OpenCL Developer Guide for Intel Processor Graphics". Und es gibt einen
Abschnitt darüber, wie man es versucht, um die überschüssige Synchronisation nicht dort abzulegen.
Aus meiner Sicht ist dies also im Prinzip eine nichtmenschliche Aufgabe. Die Leute sind extrem schlecht darin, über Multithread-Synchronisation nachzudenken, und tendieren dazu, Synchronisation entweder konservativ oder falsch oder beides auf einmal zu formen - ich habe solche Kommas gesetzt (
aber wir haben es behoben - Anmerkung der Redaktion ).
Auf der anderen Seite schreiben Sie in DPC ++ keinen Code mit expliziten Barrieren wie folgt:
for (t = 0; t < numTiles; t++) { const int tiledRow = TS * t + row; const int tiledCol = TS * t + col; Asub[col][row] = A[globalRow * AY + tiledCol]; Bsub[col][row] = B[tiledRow * BY + globalCol];
Sie werden höchstwahrscheinlich eine explizite Iteration von
parallel_for_work_group schreiben, innerhalb der
group.parallel_for_work_item cgh.parallel_for_work_group<class mxm_kernel>( cl::sycl::range<2>{BIG_AX / TS, BIG_BY / TS}, cl::sycl::range<2>{TS, TS}, [=](cl::sycl::group<2> group) {
Infolgedessen müssen Sie die Synchronisation überhaupt nicht manuell einstellen, und der gesamte Abschnitt kann weggeworfen werden.
Und so können Sie in allen Abschnitten laufen. Etwas wird überleben, etwas wird gehen. Ich sehe das Erscheinen eines neuen Dokuments „Optimierung für DPC ++“ voraus, aber die Zeit sollte vergehen, da alle wirklich funktionierenden Techniken erst später und mit Blut entwickelt werden
Frage aus der Mail
In OpenCL gibt es eine Einschränkung: Sie können keine "entfernten Daten" im Kernel verwenden, dh Sie implementieren beispielsweise einen "breiten Filter", der Eingabedaten aus einer großen Gruppe von Pixeln verwendet, die größer sind als die OpenCL-Arbeitsgruppe in einer Berechnung. Was bietet DPC ++ in dieser Hinsicht?
Die antwort. Das ist unmöglich. Natürlich schreibe ich keine besonderen Kernel ... Aber es ist absolut sicher, dass Sie den gesamten globalen Speicher so nutzen können, wie er ist. Sie müssen nur sicherstellen, dass Sie mit atomaren Operationen arbeiten (oder hierarchische Kernel extern synchronisieren). Sie können auch System-SVM (oder USM in DPC ++) anschließen.
Leider ist das alles extrem ineffizient, und ich mag all diese Tricks nicht. Außerdem sind sie vom Compiler nur schwer zu optimieren.
Wenn wir also von direkten und effektiven Lösungen sprechen, gibt es in DPC ++ natürlich keine Magie. Ihr Programm ist am Ende immer noch in Teile geteilt: Der Hostcode und der Gerätecode, und alle Geräteeinschränkungen wirken sich auf den Gerätecode aus. Die maximale Größe der Arbeitsgruppe ist die tatsächliche Parallelität, zu der Ihre Hardware fähig ist. Alles, was noch dazu kommt, sind Wege, um herauszukommen und die Leistung dramatisch zu beeinträchtigen. Aus diesem Grund bietet DPC ++ die Möglichkeit, dies zu tun:
device.get_info <sycl :: info :: device :: max_work_group_size> () und dann zu entscheiden, wie mit der resultierenden Zahl
umgegangen werden soll .
Es wäre natürlich verlockend, ein Modell in DPC ++ zu erstellen, wenn der Programmierer mit beliebig langen Schleifen arbeitet und der Compiler überlegt, was als nächstes zu tun ist, aber es wäre tödlich, weil es Konstanten und sogar Asymptoten mit zusätzlicher Komplexität verbergen würde Computer tauchen aus dem Nichts auf. Aus einem anderen Grund schrieb Alexandrescu, dass „Komplexität als Verbrechen betrachtet werden sollte“, und dies gilt auch.
Manchmal hilft es, den Algorithmus selbst zu überarbeiten. Hier erleichtert DPC ++ die Arbeit, da strukturierter Code einfacher umzugestalten ist. Aber das ist so ein Trost.
Frage aus der Mail
DPC ++ basiert auf SYCL. Aber was ist, wenn Sie tiefer unter die Haube gehen, was sind die Unterschiede zu OpenCL bei der Implementierung des Backends, falls vorhanden? Ist beispielsweise der Verteilungsmechanismus zwischen heterogenen Geräten der gleiche wie bei OpenCL?
Die antwort. Wenn Sie unter die Haube kommen, dann ist dies OpenCL. Alle Vorteile und Stärken von SYCL sind die Vorteile und Stärken der Sprache, dh des Frontends. Vom Frontend kommt die gute alte SPIRV, die zum Backend geht und dort bereits für eine bestimmte Grafikkarte optimiert ist (oft bereits zur Laufzeit, d. H. Es ist JIT), genau wie OpenCL dafür optimiert wäre.
Eine andere Sache ist, dass der Mechanismus für die Verteilung der Arbeit zwischen heterogenen Geräten mehr Front-End als Back-End ist, da es der Host-Code ist, der entscheidet, was gesendet wird und wohin. Der Hostcode wird von DPC ++ bezogen. Ich habe bereits ein etwas höheres Beispiel gezeigt, wie der Compiler auf der Basis von Accessoren eine Entscheidung über die Parallelisierung von Paketen treffen kann. Und das ist nur die Spitze des Eisbergs.
Frage aus der Mail
Bibliotheken Ja, wir sprechen nicht über CUDA. Wir wissen jedoch, dass es für CUDA-Entwickler sehr nützliche Bibliotheken gibt, die auf der GPU mit hoher Leistung arbeiten. OneAPI enthält auch einige Bibliotheken, aber zum Beispiel IPP - für die Arbeit mit Bildern in oneAPI / OpenCL ist keine Archivierung sinnvoll. Wird es etwas geben und wie kann in diesem Fall von CUDA auf oneAPI gewechselt werden?
Die antwort. Der Übergang von CUDA zu einem einzigen offenen Standard wird schwierig, aber unvermeidlich sein. Natürlich verfügt CUDA jetzt über eine ausgereiftere Infrastruktur. Die Eigenschaften der Lizenzierung sind jedoch ein Nachteil, da immer mehr Akteure auf dem Markt für heterogene Systeme, immer interessantere Karten und Beschleuniger verschiedener Hersteller auftreten.
Die Vielfalt der vorhandenen APIs erschwert es Programmierern mit Erfahrung in der klassischen CPU, diese Welt der Möglichkeiten zu nutzen. Was zu OneAPI oder so ähnlich führt. Hier liegt die Magie nicht im Durchbruch von Intel in der Grafik, sondern in der Tatsache, dass Intel jedem die Tür zu DPC ++ öffnet. Wir besitzen nicht einmal den SYCL-Standard, er gehört zur Khronos-Gruppe, und alle Intel-Erweiterungen sind Erweiterungen in Khronos, für die sich jeder engagieren kann (und es gibt Vertreter aller wichtigen Akteure). Und das bedeutet, dass (Bibliotheken) und Community angezeigt werden (bereits angezeigt werden) und eine Reihe von Stellen in diese Richtung.
Und natürlich wird IPP für neue Realitäten umgeschrieben. Ich habe nichts mit IPP zu tun, aber die Verwendung von DPC ++ ist gesunder Menschenverstand.
Aber was noch wichtiger ist, jetzt ist genau der Moment in der Geschichte, in dem Sie Ihre eigene Bibliothek schreiben können, die IPP übertrifft und die dann von der ganzen Welt verwendet wird. Weil offene Standards immer gewinnen.
Frage aus der Mail
Wenn wir den Start von Trainings- und Inferenz-Algorithmen für neuronale Netze auf Nervana und FPGA vergleichen, welche Unterschiede bestehen in der Programmierung und der daraus resultierenden Effizienz?
Die antwort. Ich weiß nichts über FPGA-Programmierdetails, ich schreibe Compiler. Aber ich habe eine Gegenfrage. Und wie werden wir vergleichen? Bei Standard-Benchmarks ist es unsportlich, Nervana leckte darunter. Aber wenn Sie etwas Interessantes haben, wird das FPGA Ihre Hände lösen, und wenn Sie Nervana dieses Etwas geben, kann das langwierig und teuer sein, das ist alles.
Es stellt sich heraus, dass die Frage selbst sozusagen aus der Serie „Wer ist stärker als ein Elefant oder ein Wal?“ Stammt. Dies ist jedoch keine wirkliche Frage. Die eigentliche Frage lautet: Wie kann man einen Elefanten und einen Wal in einem Wagen fangen? Nun, oder zumindest verteilen, sagen wir, dass ein Elefant ihn an Land zieht und ein Wal auf dem Seeweg.
Im Fall von OneAPI haben Sie im Allgemeinen dasselbe Programm in Standard-C ++. Und Sie können es selbst schreiben und mit hin und her auslagern. Dies ist genau die Aufgabe, die Sie interessiert, an der Sie selbst die Leistung messen und optimieren können. Ein einziger Standard und eine einzige Schnittstelle zu heterogenen Geräten sind ein Schritt zum Vergleich von Äpfeln mit Äpfeln in solchen Angelegenheiten.
Zum Beispiel: "Was ist für% meiner Aufgabe% im Hinblick auf einfache Programmierung und Effizienz besser? Setzen Sie diesen Teil auf FPGA, belassen Sie diesen auf Nervana oder teilen Sie diesen Teil in zwei Teile und schreiben Sie diesen Teil für die GPU neu."
Und die ganze Geschichte mit OneAPI - Sie müssen nur sagen: "Warum lange darüber nachdenken, ich werde es jetzt schnell versuchen, ES IST EINFACH."
Noch nicht so einfach. Aber es wird geben.
Nachwort vom ExpertenVielen Dank für Ihre Fragen. Es ist möglich und sogar wahrscheinlich, dass ich falsch, ungenau und fehlerhaft lag. Es passiert, im Internet irrt sich ständig jemand.
Ich hoffe, dass ich jemanden für heterogene Programmierung und DPC ++ interessieren konnte. Ich möchte allen die Seite
sycl.tech empfehlen, auf der
jede Menge Berichte liegen, auch von weltbekannten Experten (Englisch ist erforderlich).
Gut zu allen!
PS vom Verlag. Diesmal wurde durch einstimmige Entscheidung der Redaktion beschlossen, den Preis für die beste Frage ... an den Autor der Antworten zu vergeben. Ich denke, Sie werden zustimmen, dass dies fair ist.