7
Ich möchte Assembler-Code in CUDA C-Code verwenden, um teure Ausführungen wie wir mit asm in c-Programmierung zu reduzieren.Ist es möglich, Montageanweisungen in den CUDA-Code aufzunehmen?
Ist es möglich?
A
Antwort
4
Nein, Sie können nicht, es gibt nichts wie die asm Konstrukte von C/C++. Was Sie tun können, ist die generierte PTX-Baugruppe zu optimieren und dann mit CUDA zu verwenden.
Siehe this für ein Beispiel.
Aber für GPUs sind Assembly-Optimierungen NICHT notwendig, Sie sollten zuerst andere Optimierungen durchführen, wie Speicherkoaleszenz und Belegung. Weitere Informationen finden Sie unter CUDA Best Practices guide.
+2
Zweitens, dass! Meiner Erfahrung nach sind CUDA-Programme fast immer speichergebunden, nicht compute-gebunden. – mch
+0
danke über beide. Ich wollte nur die Anzahl der Divisions- und Modulo-Operationen reduzieren, aber jetzt konzentriere ich mich auf das Speicherproblem. – superscalar
+0
Hinweis, wenn Sie gegen die neueste Architektur kompilieren (mit dem Flag -arch SM_20), ist die neueste API jetzt vollständig ?? konform mit IEEE-Gleitkomma-Spezifikationen für Division und Quadratwurzel. Wenn Sie eine Reihe von Divisionen haben und Sie auch den Befehl -arch sm_20 verwenden, sollten Sie möglicherweise mit einem Flag auf die "weniger" -kompatible Version umsteigen, um die Leistung zu steigern: __- prec-div = false__ http://forums.nvidia.com/lofiversion/index.php?t170749.html –
17
Seit CUDA 4.0 wird Inline-PTX von der CUDA Toolchain unterstützt. Es gibt ein Dokument im Toolkit, das es beschreibt: Using_Inline_PTX_Assembly_In_CUDA.pdf
Unten finden Sie einen Code, der die Verwendung von Inline-PTX in CUDA 4.0 demonstriert. Beachten Sie, dass dieser Code nicht als Ersatz für die integrierte __clz() -Funktion von CUDA verwendet werden sollte. Ich habe ihn lediglich geschrieben, um Aspekte der neuen Inline-PTX-Funktion zu untersuchen.
__device__ __forceinline__ int my_clz (unsigned int x)
{
int res;
asm ("{\n"
" .reg .pred iszero, gezero;\n"
" .reg .u32 t1, t2;\n"
" mov.b32 t1, %1;\n"
" shr.u32 %0, t1, 16;\n"
" setp.eq.b32 iszero, %0, 0;\n"
" mov.b32 %0, 0;\n"
"@iszero shl.b32 t1, t1, 16;\n"
"@iszero or.b32 %0, %0, 16;\n"
" and.b32 t2, t1, 0xff000000;\n"
" setp.eq.b32 iszero, t2, 0;\n"
"@iszero shl.b32 t1, t1, 8;\n"
"@iszero or.b32 %0, %0, 8;\n"
" and.b32 t2, t1, 0xf0000000;\n"
" setp.eq.b32 iszero, t2, 0;\n"
"@iszero shl.b32 t1, t1, 4;\n"
"@iszero or.b32 %0, %0, 4;\n"
" and.b32 t2, t1, 0xc0000000;\n"
" setp.eq.b32 iszero, t2, 0;\n"
"@iszero shl.b32 t1, t1, 2;\n"
"@iszero or.b32 %0, %0, 2;\n"
" setp.ge.s32 gezero, t1, 0;\n"
" setp.eq.b32 iszero, t1, 0;\n"
"@gezero or.b32 %0, %0, 1;\n"
"@iszero add.u32 %0, %0, 1;\n\t"
"}"
: "=r"(res)
: "r"(x));
return res;
}
Verwandte auf SU: http://superuser.com/questions/668019/how-do-device-driver-instructions-program-the-gpu Für Intel Xeon Phi sieht es aus möglich. –
Mögliches Duplikat von [Wie man GPU-Assembler erstellt oder manipuliert?] (Http://stackoverflow.com/questions/4660974/how-to-create-manipulate-gpu-assembler) –