große integer-addition mit CUDA
Ich habe die Entwicklung eines kryptografischen Algorithmus auf der GPU und derzeit stecken mit einem Algorithmus durchführen zu großen integer-addition. Große Ganzzahlen dargestellt werden, in einer üblichen Weise als eine Reihe von 32-bit-Worten.
Zum Beispiel, können wir einen thread zum addieren zwei 32-bit-Worten. Für Einfachheit, lassen vermuten
dass die zahlen, die Hinzugefügt werden, sind von der gleichen Länge und die Anzahl der threads pro block == Anzahl der Wörter. Dann:
__global__ void add_kernel(int *C, const int *A, const int *B) {
int x = A[threadIdx.x];
int y = B[threadIdx.x];
int z = x + y;
int carry = (z < x);
/** do carry propagation in parallel somehow ? */
............
z = z + newcarry; //update the resulting words after carry propagation
C[threadIdx.x] = z;
}Ich bin mir ziemlich sicher, dass es einen Weg gibt, zu tragen, Ausbreitung über einige knifflige Reduzierung Verfahren aber konnte es nicht herausfinden..
Hatte ich einen Blick auf CUDA-Schub-Erweiterungen aber große integer-Paket scheint nicht zu realisieren noch.
Vielleicht kann jemand mir einen Tip geben, wie zu tun, die auf CUDA ?
- Die GPU kann mit bis zu 64 bit (long long) direkt. Ein Ansatz für 128-bit ist beschrieben im dies ALSO Frage/Antwort.
- Ich denke, was willst du von CUDA erreicht werden kann, von C-Techniken. Daher habe ich retaged die Frage in
Czu. Hoffe auf nette Antwort von den C-Experten. - Ja, Sie können auch das Programm einer long-integer-addition mit nur high-level-C-Konstrukte (im Gegensatz zu PXT linline Montage in CUDA), aber es würde deutlich mehr Anweisungen, als ich darauf hinwies, in dieser Antwort auf: stackoverflow.com/questions/12448549/...
- vielen Dank für die Anregungen. Ich weiß, dass CUDA unterstützt spezielle systeminterne Funktionen zu verwenden carry-flag nach Ergänzungen. Der Punkt ist, die zahlen können sehr groß (über 2048 32-bit-Worte) also ich bin wirklich auf der Suche für eine parallele Lösung, vielleicht über parallele Reduktion irgendwie ?
- Außerdem ist überträgt sich rein rechnerisch nicht intensiv genug ist, um sinnvoll aufgeteilt auf mehrere threads (zumindest aus der Spitze von meinem Kopf). Für die Multiplikation, Sie könnten jeden thread arbeiten auf der Summierung einer Spalte teilweise 32x32->64-bit-Produkte, dann propagieren der trägt am Ende. Sie konnten auch einen Blick in die latenten tragen Vermehrung, indem Sie die Ergebnisse einer Zusatz as separate Summen-und carry-Vektoren. Viel hängt von der genauen algorithmischen Kontext.
- ja, ich Stimme zudem nicht, dass rechenintensive als Multiplikation. Aber immer noch, wenn ich eine integer der Länge 2048 oder noch mehr Worte, dabei zusätzlich in einer Schleife mit einem CUDA-thread wäre sehr innefficient, weil dieser Vorgang sieht embarrasingly parallel zu mir, außer für die Fortpflanzung.
- Ich habe gehackt zusammen ein cuda-kernel zu tun, die parallele Zugabe von bis zu 1024 64-bit-unsigned-Mengen, und in der Lage zu handhaben Chargen dieser Probleme parallel auch. Ein kernel-Berechnung Standpunkt, und wenn wir Chargen einer großen Anzahl von Problemen in parallel, es ist über 10x schneller als meine naiven CPU-code. Wenn Sie werfen in die Daten kopieren Zeit, es ist in etwa auf Augenhöhe mit der CPU-Zeit. Es gibt keine PTX -, nur C-code, also ich bin mir sicher, dass es gemacht werden könnte, um schneller zu laufen, aber ich weiß nicht, wie viel. Ich poste es als Antwort, wenn Sie wollen, um es zu betrachten. Auch mache ich keine Ansprüche über seine Richtigkeit.
- Den folgenden link für CUDA-Schub-Erweiterungen enthält den code für große integer-Multiplikation: cuda-thrust-extensions.googlecode.com/svn/trunk/big%20integer
Du musst angemeldet sein, um einen Kommentar abzugeben.
Sind Sie richtig, tragen Vermehrung kann über prefix-sum-Berechnung, aber es ist ein bisschen schwierig zu definieren, die binäre Funktion für diesen Vorgang und beweisen, dass es ist assoziativ (wird benötigt, für die parallele Präfix-Summe). Als eine Angelegenheit von der Tat, dieser Algorithmus wird verwendet, (theoretisch) in Carry-lookahead-adder.
Angenommen, wir haben zwei große ganze zahlen a[0..n-1] und b[0..n-1].
Dann berechnen wir (i = 0..n-1):
Definieren wir zwei Funktionen:
mit Recht intuitive Bedeutung: generieren[i] == 1 bedeutet, dass der carry generiert am
position, die ich während propagieren[i] == 1 bedeutet, dass das tragen weitergegeben wird, von der position
(i - 1) (i + 1). Unser Ziel ist die Berechnung der Funktion carryout[0..n-1] verwendet zum aktualisieren der resultierende Summe s[0..n-1]. carryout berechnet werden kann, die rekursiv wie folgt:
Hier carryout[i] == 1 if carry erzeugt an der position i ODER es wird generiert, manchmal früher UND vermehrt auf position habe ich. Endlich, wir aktualisieren die resultierende Summe:
Nun es ist ziemlich einfach zu beweisen, dass carryout-Funktion ist in der Tat binären, assoziativen und daher parallele Präfix-Summe Berechnung gilt. Um dies zu implementieren, die auf CUDA, wir können mischen der beiden flags 'erzeugen' und 'verbreiten' in eine einzelne variable, da Sie sich gegenseitig ausschließen, d.h.:
In anderen Worten,
Dann kann man überprüfen, dass die folgende Formel berechnet Präfix-Summe für carryout Funktion:
für alle k < ich. Der Beispielcode unten zeigt eine große Bereicherung für 2048-word-Integer-zahlen. Hier habe ich noch CUDA-Blöcke mit 512 threads:
Beachten Sie, dass Makros UADDO /UADDC vielleicht nicht mehr notwendig, seit CUDA 4.0 verfügt über die entsprechenden Interna (aber ich bin mir nicht ganz sicher).
Auch die Bemerkung, dass, obwohl die parallele Reduktion ist Recht schnell, wenn Sie brauchen, um hinzuzufügen mehrere große zahlen in einer Zeile, es wäre besser, verwenden einige redundante Darstellung (die wurde vorgeschlagen, in den Kommentaren oben), d.h., zuerst sammeln die Ergebnisse der Additionen im 64-bit-Worten, und führen Sie dann einen carry-propagation am Ende in "einem Zug".
Ich dachte, ich würde post auf meine Antwort auch, zusätzlich zu @asm, also ist diese Frage ALSO eine Art von repository-Ideen. Ähnlich wie @asm ich Sie erkennen und speichern der carry-Zustandes sowie die "carry-through" - Zustand, dh. wenn die intermediate-Wort Ergebnis wird alle 1 s (0xF...FFF), so dass, wenn Sie einen tragen, wurden zu vermehren, in dieses Wort, es würde "carry-through", um das nächste Wort.
Ich habe keine PTX oder asm in meinem code, also wählte ich für die Verwendung von 64-bit vorzeichenlosen Ganzzahlen anstelle von 32-bit zu erreichen, 2048x32bit Fähigkeit, mit 1024 threads.
Einen größeren Unterschied aus @asm code ist in meiner parallel-carry-propagation-Schema. Ich konstruiere ein bit-packed array ("carry"), wo jedes bit repräsentiert das carry-Bedingung generiert, die von der unabhängigen intermediate 64-bit ergänzt, die aus jedem der 1024 threads. Ich auch konstruieren, eine bit-packed array ("carry_through"), wo jedes bit stellt den carry_through Zustand der einzelnen 64-bit-zwischenergebnisses. Für 1024 threads, diese Beträge zu 1024/64 = 16x64-bit-Wörtern von shared memory für jedes bit-verpackt-array, also insgesamt shared mem usage ist 64+3 32bit Mengen?. Mit diesen bit packed arrays, führe ich die folgende zu erzeugen, die eine kombinierte propagiert tragen Kennzeichen:
(beachten Sie, dass carry nach Links verschoben, indem man: carry[i] zeigt an, dass das Ergebnis von a[i-1] + b[i-1] erzeugt eine tragen)
Die Erklärung ist wie folgt:
interagieren Sie mit einer Sequenz von einem oder mehr tragen, obwohl die Bedingungen
geändert hat bits repräsentieren alle Wörter, die betroffen sein werden von
die Ausbreitung der carry in der carry_through Sequenz
zeigt den betroffenen die Ergebnisse weisen mit einem 1-bit -
tragen Indikatoren gibt eine Kombination carry-Zustand, das ist dann
update alle Zwischenergebnisse.
Beachten Sie, dass die Zugabe in Schritt 2 erfordert ein anderes multi-word hinzufügen (für große Ganzzahlen, bestehend aus mehr als 64 Worte). Ich glaube, dass dieser Algorithmus funktioniert, und es hat den test bestanden Fällen, die ich geworfen haben es.
Hier ist mein Beispiel code implementiert dies: