Je veux utiliser le code d'assemblage dans le code CUDA C afin de réduire les exécutions coûteuses comme nous le faisons en utilisant asm en programmation c.Est-il possible de mettre des instructions d'assemblage dans le code CUDA?
Est-ce possible?
Non, vous ne pouvez pas, il n'y a rien comme les constructions asm de C/C++. Ce que vous pouvez faire est de modifier l'ensemble PTX généré, puis l'utiliser avec CUDA.
Voir this pour un exemple. Mais pour les GPU, les optimisations d'assemblage ne sont PAS nécessaires, vous devez d'abord effectuer d'autres optimisations, telles que la coalescence et l'occupation de la mémoire. Voir le CUDA Best Practices guide pour plus d'informations.
Ensuite! Dans mon expérience, les programmes CUDA sont presque toujours liés à la mémoire, pas calculés liés. – mch
merci ci-dessus à la fois. Je voulais juste réduire le nombre d'opérations de division et modulo, mais maintenant je vais me concentrer sur le problème de la mémoire. – superscalar
Notez, si vous compilez contre l'architecture la plus récente (en utilisant le drapeau -arch sm_20), l'API la plus récente est maintenant entièrement ?? conforme aux spécifications de virgule flottante IEEE pour la division et la racine carrée. Si vous avez un tas de divisions et que vous utilisez aussi -arch sm_20, alors vous pourriez envisager de revenir à la version "moins" compatible pour un gain de performance en utilisant l'indicateur: __- prec-div = false__ http://forums.nvidia.com/lofiversion/index.php?t170749.html –
Depuis CUDA 4.0, PTX en ligne est pris en charge par la chaîne d'outils CUDA. Un document de la boîte à outils le décrit: Using_Inline_PTX_Assembly_In_CUDA.pdf
Voici un code montrant l'utilisation de PTX en ligne dans CUDA 4.0. Notez que ce code ne doit pas être utilisé pour remplacer la fonction __clz() intégrée de CUDA, je l'ai simplement écrit pour explorer certains aspects de la nouvelle fonctionnalité PTX intégrée.
__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;
}
connexes sur SU: http://superuser.com/questions/668019/how-do-device-driver-instructions-program-the-gpu Pour Xeon Phi Intel il semble possible. –
Copie possible de [Comment créer ou manipuler un assembleur GPU?] (Http://stackoverflow.com/questions/4660974/how-to-create-or-manipulate-gpu-assembler) –