14

Linux Ubuntu 10.04にcuda SDKをインストールできました。私のグラフィック カードは NVIDIA geForce GT 425M で、重い計算問題に使用したいと考えています。私が疑問に思っているのは、unsigned 128 ビット int var を使用する方法はありますか? gcc を使用して CPU 上でプログラムを実行する場合、__uint128_t 型を使用していましたが、cuda で使用するとうまくいかないようです。cuda で 128 ビット整数を使用するためにできることはありますか?

4

3 に答える 3

52

最高のパフォーマンスを得るには、uint4 などの適切な CUDA ベクトル型の上に 128 ビット型をマップし、PTX インライン アセンブリを使用して機能を実装します。追加は次のようになります。

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;
}

乗算は、PTX インライン アセンブリを使用して 128 ビットの数値を 32 ビットのチャンクに分割し、64 ビットの部分積を計算して適切に加算することにより、同様に構築できます。明らかに、これには少し手間がかかります。数値を 64 ビットのチャンクに分割し、__umul64hi() を通常の 64 ビット乗算およびいくつかの加算と組み合わせて使用​​することにより、C レベルで妥当なパフォーマンスを得ることができます。これにより、次のようになります。

__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;
}

以下は、PTX インライン アセンブリを使用する 128 ビット乗算のバージョンです。CUDA 4.2 に同梱されている PTX 3.0 が必要であり、コードには少なくとも計算能力 2.0 を備えた GPU、つまり Fermi または Kepler クラスのデバイスが必要です。128 ビットの乗算を実装するには 16 回の 32 ビット乗算が必要なため、このコードでは最小数の命令を使用しています。比較すると、CUDA 組み込み関数を使用する上記のバリアントは、sm_20 ターゲットに対して 23 命令にコンパイルされます。

__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;
}
于 2011-06-02T21:16:44.870 に答える
13

CUDA は 128 ビット整数をネイティブにサポートしていません。2 つの 64 ビット整数を使用して、自分で操作を偽造できます。

この投稿を見てください：

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;
}
于 2011-05-28T15:28:44.770 に答える
3

かなり遅れた答えですが、このライブラリの使用を検討できます。

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

これは、128 ビット サイズの構造体を定義し、メソッドと独立したユーティリティ関数を使用して期待どおりに機能させ、通常の整数のように使用できるようにします。多くの場合。

于 2018-05-30T13:00:17.577 に答える