7
Chcę użyć kodu zespołu w kodzie CUDA C w celu zmniejszenia kosztownych wykonań , jak to robimy przy użyciu asm w programowaniu c.Czy można umieścić instrukcje montażu w kodzie CUDA?
Czy to możliwe?
A
Odpowiedz
4
Nie, nie można, nie ma nic takiego jak konstrukty asm z C/C++. Możesz zmodyfikować wygenerowany zespół PTX, a następnie użyć go z CUDA.
Zobacz przykład na this.
Jednak w przypadku procesorów graficznych optymalizacje zespołów NIE są konieczne, należy najpierw przeprowadzić inne optymalizacje, takie jak koalescencja pamięci i obłożenie. Więcej informacji można znaleźć w CUDA Best Practices guide.
+2
Po drugie! Z mojego doświadczenia wynika, że programy CUDA prawie zawsze są związane z pamięcią, a nie są obliczane. – mch
+0
dzięki powyżej. Chciałem tylko zmniejszyć liczbę operacji dzielenia i modulo, ale teraz skupię się na problemie z pamięcią. – superscalar
+0
Uwaga, jeśli kompilujesz się z najnowszą architekturą (używając flagi -arch sm_20), najnowszy interfejs API jest teraz w pełni? zgodne ze specyfikacjami zmiennoprzecinkowymi IEEE dla dzielenia i pierwiastka kwadratowego. Jeśli masz kilka dywizji i używasz również -archa sm_20, możesz rozważyć powrót do "mniej" zgodnej wersji dla zwiększenia wydajności za pomocą flagi: __- prec-div = false__ http://forums.nvidia.com/lofiversion/index.php?t170749.html –
17
Od wersji CUDA 4.0, wbudowana funkcja PTX jest obsługiwana przez pakiet narzędziowy CUDA. W zestawie narzędzi opisanym w dokumencie znajduje się dokument: Using_Inline_PTX_Assembly_In_CUDA.pdf
Poniżej znajduje się kod demonstrujący użycie wbudowanego PTX w CUDA 4.0. Zauważ, że ten kod nie powinien być użyty jako zamiennik wbudowanej funkcji CUDA __clz(), napisałem ją tylko po to, by zbadać aspekty nowej wbudowanej funkcji PTX.
__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;
}
powiązanymi na SU: http://superuser.com/questions/668019/how-do-device-driver-instructions-program-the-gpu Dla Intel Xeon Phi wygląda to możliwe. –
Możliwy duplikat [Jak utworzyć lub manipulować asemblerem GPU?] (Http://stackoverflow.com/questions/4660974/how-to-create-or-manipulate-gpu-assembler) –