Bearbeiten: Diese Frage ist eine re-date-Version des Originals, daher sind die ersten Antworten möglicherweise nicht mehr relevant.
Ich bin gespannt, welche Auswirkungen eine Gerätefunktion mit erzwungenem No-Inlining auf die Synchronisation innerhalb einer Gerätefunktion hat. Ich habe einen einfachen Testkern, der das betreffende Verhalten veranschaulicht.
Der Kernel nimmt einen Puffer und übergibt ihn zusammen mit einem gemeinsamen Puffer und einer Indikatorvariablen, die einen einzelnen Thread als "Boss" -Thread identifiziert, an eine Gerätefunktion. Die Gerätefunktion hat einen abweichenden Code: Der Chef-Thread verbringt zuerst Zeit mit trivialen Operationen auf dem gemeinsam genutzten Puffer und schreibt dann in den globalen Puffer. Nach einem Synchronisierungsaufruf schreiben alle Threads in den globalen Puffer. Nach dem Kernel-Aufruf gibt der Host den Inhalt des globalen Puffers aus. Hier ist der Code:
CUDA-CODE:
test_main.cu
%Vor%test_kernel.cu
%Vor%Ich habe diesen Code aus dem CUDA SDK kompiliert, um die Funktionen von "cutilsacall ()" in test_main.cu zu nutzen, aber natürlich können diese auch entfernt werden, wenn Sie außerhalb des SDK kompilieren möchten. Ich kompilierte mit CUDA Driver / Toolkit Version 4.0, Rechenkapazität 2.0, und der Code wurde auf einer GeForce GTX 480 ausgeführt, die die Fermi-Architektur hat.
Die erwartete Ausgabe ist
0 1 2 3 ... blockDim.x-1
Die Ausgabe, die ich bekomme, ist jedoch
1 1 2 3 ... blockDim.x-1
Dies deutet darauf hin, dass der Boss-Thread den bedingten "scratchBuffer [0] = 1;" Anweisung NACH allen Threads führen die "scratchBuffer [threadIdx.x] = threadIdx.x;" Anweisung, obwohl sie durch eine __syncthreads () - Barriere getrennt sind.
Dies tritt auf, selbst wenn der Boss-Thread angewiesen wird, einen Sentinel-Wert in die Pufferposition eines Threads in seinem gleichen Warp zu schreiben; Der Sentinel ist der letzte im Puffer vorhandene Wert und nicht die entsprechende threadIdx.x.
Eine Änderung, die bewirkt, dass der Code die erwartete Ausgabe erzeugt, besteht darin, die bedingte Anweisung zu ändern
if (isBoss) {
bis
if (IS_BOSS ()) {
; h., um zu ändern, dass die Divergenzsteuervariable in einem Parameterregister gespeichert wird, um in einer Makrofunktion berechnet zu werden. (Beachten Sie die Kommentare zu den entsprechenden Zeilen im Quellcode.) Auf diese spezielle Änderung habe ich mich konzentriert, um das Problem aufzuspüren. Betrachtet man die disassemblierten .cubins des Kernels mit dem 'isBoss'-bedingten (dh gebrochenen Code) und dem' IS_BOSS () 'bedingten (dh Arbeitscode), so scheint der auffälligste Unterschied in den Anweisungen das Fehlen von zu sein ein SSY-Befehl im zerlegten Code.
Hier sind die disassemblierten Kernel, die durch das Zerlegen der .cubin-Dateien mit erstellt wurden "cuobjdump -sass test_kernel.cubin". alles bis zum ersten 'EXIT' ist der Kernel und alles danach ist die Gerätefunktion. Die einzigen Unterschiede sind in der Gerätefunktion.
ZERLEGTER OBJEKTCODE:
"gebrochener" Code
%Vor%"Arbeitscode"
%Vor%Die Anweisung "SSY" ist im Arbeitscode vorhanden, aber nicht im gebrochenen Code. Das Handbuch zu cuobjdump beschreibt die Anweisung mit "Synchronisationspunkt festlegen; wird vor potenziell divergierenden Anweisungen verwendet". Das lässt mich denken, dass der Compiler aus irgendeinem Grund die Möglichkeit der Divergenz im gebrochenen Code nicht erkennt.
Ich fand auch, dass, wenn ich die __noinline__ Direktive auskommentiere, der Code die erwartete Ausgabe erzeugt, und in der Tat ist die Assembly, die von den ansonsten "kaputten" und "funktionierenden" Versionen erzeugt wird, exakt identisch. Das lässt mich denken, dass, wenn eine Variable über den Call-Stack übergeben wird, diese Variable nicht zur Steuerung der Divergenz und eines nachfolgenden Synchronisierungsaufrufs verwendet werden kann; der Compiler scheint in diesem Fall die Möglichkeit der Divergenz nicht zu erkennen und fügt daher keine "SSY" -Anweisung ein. Weiß jemand, ob dies tatsächlich eine legitime Einschränkung von CUDA ist, und wenn ja, ob dies irgendwo dokumentiert ist?
Vielen Dank im Voraus.