128 bits entier sur cuda?

J'ai juste réussi à installer mon cuda SDK sous Linux Ubuntu 10.04. Ma carte graphique est une NVIDIA geForce GT 425M, et j'aimerais l'utiliser pour certains gros problème de calcul.
Ce que je me demande est: est-il possible d'utiliser certains unsigned 128 bits int var? Lors de l'utilisation de gcc pour exécuter mon programme sur le PROCESSEUR, j'ai été en utilisant l' __uint128_t type, mais de l'utiliser avec cuda ne semble pas fonctionner.
Est-ce que je peux faire pour avoir de 128 bits entiers sur cuda?

Merci beaucoup
Matteo Monti
Msoft Programmation

InformationsquelleAutor Matteo Monti | 2011-05-28
  1. 49

    Pour de meilleures performances, on voudrait carte de 128 bits type sur le dessus de CUDA type de vecteur, comme uint4, et de mettre en œuvre la fonctionnalité en utilisant PTX assembly en ligne. Le plus ressemblerait à quelque chose comme ceci:

    typedef uint4 my_uint128_t;
__device__ my_uint128_t add_uint128 (my_uint128_t addend, my_uint128_t augend)
{
    my_uint128_t res;
    asm ("add.cc.u32      %0, %4, %8;\n\t"
         "addc.cc.u32     %1, %5, %9;\n\t"
         "addc.cc.u32     %2, %6, %10;\n\t"
         "addc.u32        %3, %7, %11;\n\t"
         : "=r"(res.x), "=r"(res.y), "=r"(res.z), "=r"(res.w)
         : "r"(addend.x), "r"(addend.y), "r"(addend.z), "r"(addend.w),
           "r"(augend.x), "r"(augend.y), "r"(augend.z), "r"(augend.w));
    return res;
}

    La multiplication peut également être construit à l'aide de PTX assembly en ligne par la rupture de la 128 bits en 32 bits morceaux, le calcul de l'64 bits partielle des produits et de l'ajout de façon appropriée. Évidemment, cela prend un peu de travail. On peut obtenir des performances acceptables au niveau C en brisant le nombre en 64 bits morceaux et de l'aide __umul64hi() en conjonction avec d'régulière 64 bits de multiplication et de quelques ajouts. Ce résultat suivant:

    __device__ my_uint128_t mul_uint128 (my_uint128_t multiplicand, 
                                     my_uint128_t multiplier)
{
    my_uint128_t res;
    unsigned long long ahi, alo, bhi, blo, phi, plo;
    alo = ((unsigned long long)multiplicand.y << 32) | multiplicand.x;
    ahi = ((unsigned long long)multiplicand.w << 32) | multiplicand.z;
    blo = ((unsigned long long)multiplier.y << 32) | multiplier.x;
    bhi = ((unsigned long long)multiplier.w << 32) | multiplier.z;
    plo = alo * blo;
    phi = __umul64hi (alo, blo) + alo * bhi + ahi * blo;
    res.x = (unsigned int)(plo & 0xffffffff);
    res.y = (unsigned int)(plo >> 32);
    res.z = (unsigned int)(phi & 0xffffffff);
    res.w = (unsigned int)(phi >> 32);
    return res;
}

    Ci-dessous est une version de la 128 bits de multiplication qui utilise PTX assembly en ligne. Il nécessite PTX 3.0, qui est fourni avec CUDA 4.2, et le code exige un GPU avec au moins calculer la capacité de 2.0, c'est à dire un Fermi ou Kepler de la classe de l'appareil. Le code utilise le nombre minimal d'instructions, comme seize 32 bits multiplie sont nécessaires pour mettre en œuvre un cryptage de 128 bits de multiplication. Par comparaison, la variante ci-dessus à l'aide de CUDA intrinsèques compile à 23 instructions pour une sm_20 cible.

    __device__ my_uint128_t mul_uint128 (my_uint128_t a, my_uint128_t b)
{
    my_uint128_t res;
    asm ("{\n\t"
         "mul.lo.u32      %0, %4, %8;    \n\t"
         "mul.hi.u32      %1, %4, %8;    \n\t"
         "mad.lo.cc.u32   %1, %4, %9, %1;\n\t"
         "madc.hi.u32     %2, %4, %9,  0;\n\t"
         "mad.lo.cc.u32   %1, %5, %8, %1;\n\t"
         "madc.hi.cc.u32  %2, %5, %8, %2;\n\t"
         "madc.hi.u32     %3, %4,%10,  0;\n\t"
         "mad.lo.cc.u32   %2, %4,%10, %2;\n\t"
         "madc.hi.u32     %3, %5, %9, %3;\n\t"
         "mad.lo.cc.u32   %2, %5, %9, %2;\n\t"
         "madc.hi.u32     %3, %6, %8, %3;\n\t"
         "mad.lo.cc.u32   %2, %6, %8, %2;\n\t"
         "madc.lo.u32     %3, %4,%11, %3;\n\t"
         "mad.lo.u32      %3, %5,%10, %3;\n\t"
         "mad.lo.u32      %3, %6, %9, %3;\n\t"
         "mad.lo.u32      %3, %7, %8, %3;\n\t"
         "}"
         : "=r"(res.x), "=r"(res.y), "=r"(res.z), "=r"(res.w)
         : "r"(a.x), "r"(a.y), "r"(a.z), "r"(a.w),
           "r"(b.x), "r"(b.y), "r"(b.z), "r"(b.w));
    return res;
}
    • Merci BEAUCOUP!! C'est exactement ce dont j'avais besoin!!
    • Je suppose aujourd'hui vous proposer une solution sur la base de 2 les valeurs de 64 bits?
    • Peu probable, puisque 64-bit integer opérations sont émulées et il est généralement préférable de construire des émulations sur le dessus de native instructions plutôt que d'autres émulations. Parce que entier de 32 bits de se multiplier et de se multiplier-ajouter sont eux-mêmes des émules sur Maxwell et Pascal architectures, il serait peut-être préférable d'utiliser natif 16 bits où il se multiplie de la carte à la machine instruction XMAD (16x16+32 bit multiply-add opération). J' lire que les indiens d'entier de 32 bits multiplie ont été restaurés avec de la Volta architecture , mais je n'ai pas de mains sur l'expérience avec Volta encore.
    • Comment la performance est comparée à 32 bits entiers? 1/16 ou similaires?
    • Basé sur l'instruction count il serait autour de 1/16 d'un natif 32-bit de multiplication. Le réel impact sur les performances peuvent varier un peu en fonction du code de contexte basée sur le chargement des unités fonctionnelles et registre d'utilisation.
    InformationsquelleAutor njuffa
  2. 12

    CUDA ne prend pas en charge 128 bits entiers en mode natif. Vous pouvez le faux opérations vous-même à l'aide de deux entiers 64 bits.

    Regarder ce post:

    typedef struct {
  unsigned long long int lo;
  unsigned long long int hi;
} my_uint128;

my_uint128 add_uint128 (my_uint128 a, my_uint128 b)
{
  my_uint128 res;
  res.lo = a.lo + b.lo;
  res.hi = a.hi + b.hi + (res.lo < a.lo);
  return res;
}
    • Merci beaucoup! Juste une question: à partir d'un point de vue, est-ce que ça va être assez rapide?
    • J'ai testé ce code sur mon CPU. Il fonctionne réellement, mais c'est 6 fois plus lent que d'utiliser l' __uint128_type t... n'est-ce pas là un moyen de le rendre plus rapide?
    • Vous avez testé intégrée de 128 bits entiers sur CPU avec ce my_uint128 sur le CPU? Bien sûr, la prise en charge native sera plus rapide. L'espoir est que les performances sur le GPU avec cette 128 bits type sera plus rapide que la performance sur le CPU avec construit-dans 128 bits entiers.
    InformationsquelleAutor tkerwin
  3. 2

    Une réponse tardive, mais vous pouvez envisager d'utiliser cette bibliothèque:

    https://github.com/curtisseizert/CUDA-uint128

    qui définit un cryptage de 128 bits, la taille de la structure, avec des méthodes et autonome les fonctions de l'utilitaire pour l'obtenir à fonction comme prévu, ce qui lui permet d'être utilisé comme un entier. La plupart du temps.

    • C'est vraiment cool, et bien meilleure réponse que les autres 🙂 Après avoir regardé le code source, j'ai vu qu'il y a un __mul64hi PTX instruction qui rend 64 * 64 bits multiplication efficace.
    InformationsquelleAutor einpoklum